개요
Ray Tracing in One Weekend 2장(Vec3 / Color)에서 사용하는 Vec3(Vector3) 클래스를 CUDA 디바이스 코드에서도 호출 가능하도록 포팅하고, 렌더링 커널에서 프레임버퍼 타입을 Vec3 기반으로 변경한 작업을 정리한다.
이번 변경의 핵심은 아래 두 가지다.
•
Vec3의 모든 연산이 CPU(host) / GPU(device) 양쪽에서 컴파일되도록 정리
•
커널에서 픽셀 값을 float* 3채널 배열이 아니라 Color(Vec3) 객체로 다루도록 구조 개선
변경 파일 목록
파일 | 변경 유형 | 설명 |
Vec3.h | 수정 | CUDA __host__ __device__ 지정자 추가, std:: 수학 함수 제거 |
수정 | Vec3(Color) 기반 프레임버퍼 및 렌더 커널로 변경 | |
RayTracinginOneWeekend.vcxproj | 수정 | Post-Build Event로 실행 및 output.ppm 복사 자동화 |
1. Vec3.h: CUDA 호환 포팅
1.1 변경 목적
기존 Vec3는 CPU 전용 구현이었다. CUDA 커널(__global__) 내부에서 벡터 연산을 호출하려면, 해당 함수들이 디바이스 코드로도 컴파일되어야 한다.
CUDA에서는 이를 위해 함수 선언에 다음 지정자를 붙인다.
•
__host__: CPU에서 호출 가능
•
__device__: GPU에서 호출 가능
둘 다 붙이면 “한 함수 정의로 CPU/GPU 양쪽에서 사용”이 가능해진다.
1.2 주요 변경 사항
생성자, 접근자, 길이 계산, 연산자 오버로딩 등 커널에서 직접/간접 호출될 수 있는 모든 함수에 __host__ __device__를 붙였다.
// Before (CPU 전용)
class Vec3 {
public:
Vec3() : e{0, 0, 0} {}
double x() const { return e[0]; }
// ...
};
inline Vec3 operator+(const Vec3& u, const Vec3& v) {
return Vec3(u.e[0] + v.e[0], u.e[1] + v.e[1], u.e[2] + v.e[2]);
}
C++
복사
// After (CPU + GPU)
class Vec3 {
public:
__host__ __device__ Vec3() : e{0, 0, 0} {}
__host__ __device__ double x() const { return e[0]; }
// ...
};
__host__ __device__ inline Vec3 operator+(const Vec3& u, const Vec3& v) {
return Vec3(u.e[0] + v.e[0], u.e[1] + v.e[1], u.e[2] + v.e[2]);
}
C++
복사
__global__, __host__, __device__는 키워드가 아니라 CUDA 전용 함수 지정자다.
•
__global__: 커널(호스트에서 런치)
•
__device__: 디바이스에서만 호출
•
__host__ __device__: 한 정의로 양쪽 지원
디바이스 코드에서는 std::sqrt, std::fabs, std::fmin 같은 표준 라이브러리 함수가 그대로 동작하지 않는 경우가 많다.
따라서 CUDA에서 제공하는 내장 수학 함수로 교체했다.
변경 전 | 변경 후 | 비고 |
std::sqrt | sqrt | 디바이스에서 사용 가능한 내장 함수 |
std::fabs | fabs | 디바이스 내장 함수 |
std::fmin | fmin | 디바이스 내장 함수 |
Color.h에서 별도로 정의하던 Color 타입을 Vec3.h로 통합했다.
using Point3 = Vec3;
using Color = Vec3;
C++
복사
이후 장(카메라, 재질, 광선-물체 교차, 감마 보정 등)에서 Color를 벡터로 계속 다루게 되므로, 초기에 구조를 단순화해 두면 코드가 안정적이다.
RandomDouble()이 std::mt19937 기반 CPU RNG를 사용하므로 디바이스에서 호출할 수 없다.
따라서 아래 함수들을 제거했다.
•
Vec3::Random()
•
Vec3::Random(min, max)
•
RandomInUnitDisk()
•
RandomUnitVector()
•
RandomOnHemisphere()
향후에는 cuRAND 또는 커스텀 디바이스 RNG로 대체할 예정이다.
•
사칙연산(+ - * /), 스칼라/벡터 연산
•
Dot, Cross
•
UnitVector
•
Reflect, Refract
•
Length, LengthSquared, NearZero
2. kernel.cu: Vec3 기반 렌더 커널
2.1 변경 목적
1장에서는 프레임버퍼를 float*로 두고 RGB 채널을 직접 인덱싱했다.
2장 이후에는 색/벡터 연산이 자주 등장하므로, 프레임버퍼를 Color(Vec3) 배열로 바꾸면 코드가 단순해지고 확장에 유리하다.
2.2 주요 변경 사항
// Before: float 3채널을 한 배열에 저장
float* frameBuffer;
size_t frameBufferSize = 3 * numPixels * sizeof(float);
C++
복사
// After: 픽셀 1개 = Color(Vec3) 1개
Color* frameBuffer;
size_t frameBufferSize = numPixels * sizeof(Color);
C++
복사
// Before
__global__ void render(float* frameBuffer, int maxX, int maxY) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i >= maxX || j >= maxY) return;
int pixelIndex = j * maxX * 3 + i * 3;
frameBuffer[pixelIndex + 0] = float(i) / float(maxX);
frameBuffer[pixelIndex + 1] = float(j) / float(maxY);
frameBuffer[pixelIndex + 2] = 0.2f;
}
C++
복사
// After
__global__ void render(Color* frameBuffer, int maxX, int maxY) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i >= maxX || j >= maxY) return;
int pixelIndex = j * maxX + i;
frameBuffer[pixelIndex] = Color(
double(i) / double(maxX),
double(j) / double(maxY),
0.2
);
}
C++
복사
Color col = frameBuffer[pixelIndex];
int ir = int(255.99 * col.x());
int ig = int(255.99 * col.y());
int ib = int(255.99 * col.z());
C++
복사
3. vcxproj: Post-Build Event로 실행 자동화
3.1 변경 목적
빌드(Ctrl+B) 또는 실행(F5) 시 자동으로 레이트레이서를 실행하고, 생성된 output.ppm을 솔루션 루트로 복사해서 확인 흐름을 단순화한다.
3.2 추가된 설정
Debug / Release 각각의 ItemDefinitionGroup에 동일하게 추가:
<PostBuildEvent>
<Command>
"$(OutDir)$(TargetName)$(TargetExt)"
copy /Y "$(OutDir)output.ppm" "$(SolutionDir)output.ppm"
</Command>
<Message>Running ray tracer and copying output.ppm to solution directory...</Message>
</PostBuildEvent>
XML
복사
3.3 동작 흐름
•
빌드 완료
•
생성된 exe 자동 실행
•
output.ppm 생성
•
솔루션 루트로 복사
경로 | 설명 |
x64/Debug/output.ppm | exe 실행 위치에 생성 |
(솔루션 루트)/output.ppm | Post-Build에서 복사된 파일 |
렌더링 결과(현재 단계)
•
해상도: 1440 × 720
•
블록 크기: 8 × 8 (64 threads/block)
•
그리드 크기: 181 × 91 blocks
•
출력 형식: PPM (P3, ASCII)
•
이미지 내용: R = x 비율, G = y 비율, B = 0.2 고정 그라디언트