Image render with CUDA - CUDA 에서 렌더링 하기

개요

이 글은 CUDA로 이미지(프레임버퍼)를 렌더링하는 가장 기본적인 아키텍처를 정리한다.

핵심 흐름은 아래 4단계다.

프레임버퍼를 GPU에서 접근 가능한 메모리로 준비한다.

커널을 실행해서 픽셀 단위로 병렬 색상 계산을 수행한다.

cudaDeviceSynchronize()로 GPU 작업 완료를 보장한다.

CPU에서 프레임버퍼를 읽어 PPM 파일로 저장한다.

1. CUDA 렌더링 아키텍처

1.1 CPU vs GPU 렌더링 비교

CPU(순차 처리)는 보통 “한 번에 1픽셀”을 계산한다.

for (int j = 0; j < height; ++j) {
	for (int i = 0; i < width; ++i) {
		pixel[j][i] = ComputeColor(i, j);
	}
}
C++
복사

GPU(병렬 처리)는 “수천~수백만 픽셀”을 스레드로 쪼개 동시에 계산한다.

render<<<blocks, threads>>>(frameBuffer, width, height);
C++
복사

1.2 CUDA 스레드 구조(Grid / Block / Thread)

•

Grid: 전체 작업(여기서는 이미지 전체)

•

Block: 스레드 묶음(예: 8×8 = 64 threads)

•

Thread: 픽셀 하나(혹은 몇 픽셀) 담당

Grid (전체 이미지)
├── Block(0,0)
│    ├── Thread(0,0) -> Pixel(0,0)
│    ├── Thread(1,0) -> Pixel(1,0)
│    └── ...
├── Block(1,0)
│    ├── Thread(0,0) -> Pixel(8,0)
│    └── ...
└── ...
Plain Text
복사

설정	값	설명
이미지 해상도	1440 × 720	총 1,036,800 픽셀
블록 크기	8 × 8	블록당 64 스레드
그리드 크기	181 × 91	이미지를 덮기 위한 블록 수(올림 처리 포함)

2. 핵심 코드 설명

2.1 CUDA 에러 체크 매크로

CUDA API 호출은 실패할 수 있으므로, 모든 호출을 감싸서 에러 코드와 발생 위치(파일/라인)를 출력하는 패턴이 유용하다.

// checkCuda.h (예시)
#pragma once

#include <cuda_runtime.h>
#include <iostream>
#include <cstdlib>

inline void checkCuda(cudaError_t result, const char* func, const char* file, int line)
{
	if (result != cudaSuccess) {
		std::cerr << "CUDA error = " << static_cast<unsigned int>(result)
				  << " (" << cudaGetErrorString(result) << ") at "
				  << file << ":" << line << " '" << func << "'\n";
		cudaDeviceReset();
		std::exit(99);
	}
}

#define checkCudaErrors(val) checkCuda((val), #val, __FILE__, __LINE__)
C++
복사

cudaGetErrorString(result)까지 출력하면 디버깅 속도가 훨씬 빨라진다.

2.2 렌더 커널(픽셀 병렬 처리)

각 스레드가 담당할 픽셀 좌표 (i, j)를 계산하고, 이미지 범위를 벗어나는 스레드는 즉시 종료한다.

// kernel.cu (예시)
__global__ void render(float* frameBuffer, int maxX, int maxY)
{
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	int j = threadIdx.y + blockIdx.y * blockDim.y;
	if (i >= maxX || j >= maxY) return;

	int pixelIndex = j * maxX * 3 + i * 3;
	frameBuffer[pixelIndex + 0] = float(i) / float(maxX); // R
	frameBuffer[pixelIndex + 1] = float(j) / float(maxY); // G
	frameBuffer[pixelIndex + 2] = 0.2f;                   // B
}
C++
복사

좌표	수식	설명
i (x)	threadIdx.x + blockIdx.x * blockDim.x	블록 내 스레드 위치 + 블록 오프셋
j (y)	threadIdx.y + blockIdx.y * blockDim.y	블록 내 스레드 위치 + 블록 오프셋

2.3 커널 실행 구성(Execution Configuration)

이미지가 블록 크기로 딱 나누어 떨어지지 않는 경우가 대부분이므로, 그리드 크기는 올림(ceil) 방식으로 계산한다.

const int imageWidth  = 1440;
const int imageHeight = 720;

const int blockWidth  = 8;
const int blockHeight = 8;

dim3 threads(blockWidth, blockHeight);
dim3 blocks(
	(imageWidth  + blockWidth  - 1) / blockWidth,
	(imageHeight + blockHeight - 1) / blockHeight
);

render<<<blocks, threads>>>(frameBuffer, imageWidth, imageHeight);
checkCudaErrors(cudaGetLastError());
checkCudaErrors(cudaDeviceSynchronize());
C++
복사

커널 런치 직후에는 비동기라서 에러가 늦게 드러난다.

cudaGetLastError() + cudaDeviceSynchronize() 조합으로 런치 에러/실행 에러를 즉시 잡는 습관이 좋다.

3. GPU 메모리 관리

3.1 Unified Memory(cudaMallocManaged)

학습/프로토타이핑 단계에서는 Unified Memory가 매우 편하다.

float* frameBuffer = nullptr;
size_t frameBufferSize = 3 * size_t(imageWidth) * size_t(imageHeight) * sizeof(float);
checkCudaErrors(cudaMallocManaged((void**)&frameBuffer, frameBufferSize));
C++
복사

메모리 방식	설명
cudaMalloc • cudaMemcpy	GPU 전용 할당. CPUGPU 복사를 직접 관리해야 한다.
cudaMallocManaged	CPU/GPU가 같은 포인터로 접근. 데이터 이동을 런타임이 처리한다.

3.2 전체 메모리/실행 흐름

•

cudaMallocManaged로 프레임버퍼 할당

•

render<<<...>>>로 GPU가 픽셀 계산

•

cudaDeviceSynchronize로 완료 대기

•

CPU에서 프레임버퍼 읽어 파일 저장

•

cudaFree로 해제

4. PPM 이미지 출력

4.1 PPM(P3) 포맷

PPM은 헤더 3줄 이후에 픽셀 RGB를 텍스트로 나열한다.

Plain Text
복사

4.2 파일 출력 코드(C++)

아래 코드는 float(0~1) 값을 int(0~255)로 변환해 출력한다.

#include <fstream>

std::ofstream outFile("output.ppm");
outFile << "P3\n" << imageWidth << " " << imageHeight << "\n255\n";

for (int j = imageHeight - 1; j >= 0; --j) {
	for (int i = 0; i < imageWidth; ++i) {
		size_t pixelIndex = size_t(j) * size_t(imageWidth) * 3 + size_t(i) * 3;
		int ir = int(255.99f * frameBuffer[pixelIndex + 0]);
		int ig = int(255.99f * frameBuffer[pixelIndex + 1]);
		int ib = int(255.99f * frameBuffer[pixelIndex + 2]);
		outFile << ir << " " << ig << " " << ib << "\n";
	}
}
C++
복사

레이 트레이싱 예제에서는 보통 “좌하단이 (0,0)”인 좌표계를 가정한다.

반면, 이미지 포맷/뷰어는 “좌상단이 (0,0)”처럼 보이는 경우가 많아 상하가 뒤집혀 보일 수 있다.

그래서 출력에서 j를 역순으로 내려가며 기록해 최종 이미지가 기대한 방향으로 보이게 한다.

5. Post-Build Event(자동 실행)

빌드 완료 후 exe를 자동 실행하고, 생성된 output.ppm을 솔루션 루트로 복사하도록 설정했다.

<PostBuildEvent>
	<Command>
		"$(OutDir)$(TargetName)$(TargetExt)"
		copy /Y "$(OutDir)output.ppm" "$(SolutionDir)output.ppm"
	</Command>
</PostBuildEvent>
XML
복사

렌더링 결과

•

해상도: 1440 × 720

•

이미지 내용: R = x 비율, G = y 비율, B = 0.2 그라디언트

•

출력 파일: output.ppm