개요
Ray Tracing in One Weekend 6장(안티 앨리어싱)을 CUDA GPU로 구현했다.
픽셀당 1개의 레이만 쏘면 구체 가장자리에서 “맞는다/안맞는다”가 급격히 바뀌어 계단 현상(aliasing)이 생긴다. 이를 해결하기 위해 픽셀 내부에서 여러 번 랜덤 샘플링을 수행한 뒤 평균을 내는 멀티샘플링(슈퍼 샘플링)을 적용했다.
또한 기존 CPU용 Camera에서 GPU 비호환 요소를 제거하고, GPU 전용으로 단순화한 Camera를 사용했다.
변경 파일 목록
1. 안티 앨리어싱 원리
1.1 문제: 앨리어싱(계단 현상)
샘플링 없이 픽셀 중앙에 레이 1개만 쏘면, 경계(예: 구체 가장자리)에서 픽셀이 “0 아니면 1”처럼 이분화되어 계단 모양이 나타난다.
1.2 해결: 멀티샘플링(슈퍼샘플링)
픽셀 내부의 랜덤 위치로 여러 개 레이를 쏘고 평균을 내면 경계가 부드럽게 블렌딩된다.
샘플 1개: [●] -> 중앙만 판정
샘플 N개: [·∙·] -> 픽셀 내부 여러 지점 판정 후 평균
Plain Text
복사
엄밀히는 ‘지터(jitter)된 슈퍼샘플링’에 가깝다.
랜덤 샘플이 많아질수록 계단 현상은 줄지만, 대신 연산량(=레이 수)이 늘어 렌더 시간이 증가한다.
2. cuRAND: GPU 난수 생성
2.1 왜 cuRAND인가?
CPU의 std::mt19937, rand() 같은 RNG는 GPU 커널에서 사용할 수 없다.
CUDA는 디바이스에서 사용할 수 있는 난수 라이브러리 cuRAND를 제공하므로, 픽셀 샘플링용 랜덤 오프셋에 이를 사용한다.
2.2 cuRAND 상태(curandState) 초기화
픽셀마다 독립적인 난수열을 갖도록 pixelIndex를 sequence로 사용한다.
#include <curand_kernel.h>
__global__ void renderInit(int maxX, int maxY, curandState* randState)
{
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 + i;
curand_init(
1984, // seed
pixelIndex, // sequence (픽셀마다 다르게)
0, // offset
&randState[pixelIndex]
);
}
C++
복사
파라미터 | 값 | 설명 |
seed | 1984 | 전체 시드(같아도 됨) |
sequence | pixelIndex | 픽셀마다 다른 난수열 |
offset | 0 | 시퀀스 내 시작 위치 |
2.3 난수 사용(픽셀 내부 오프셋)
curand_uniform()은 (0, 1] 범위의 균등 분포 값을 준다.
curandState local = randState[pixelIndex];
float rx = curand_uniform(&local);
float ry = curand_uniform(&local);
// ... 사용 후에는 상태를 다시 저장하는 게 안전하다.
randState[pixelIndex] = local;
C++
복사
커널에서 curandState를 로컬 변수로 복사해서 쓰는 이유는 레지스터에서 계산하게 해 성능이 좋고 코드가 단순해지기 때문이다.
대신 커널이 끝나기 전에 갱신된 상태를 randState[pixelIndex] = local;로 다시 저장해야 다음 프레임/다음 커널 호출에서 난수열이 이어진다.
3. Camera.h: GPU 전용 카메라로 간소화
3.1 변경 목적
기존 CPU 버전 카메라는 std::tan, CPU RNG 등 GPU에서 바로 쓰기 어려운 요소가 섞이기 쉽다.
이 단계에서는 뷰포트를 고정한 “최소 카메라”로 단순화해서 GPU에서 바로 Ray를 생성하도록 했다.
3.2 구현 예시(C++/CUDA C++)
// Camera.h (예시)
#pragma once
#include "Ray.h"
#include "Vec3.h"
class Camera {
public:
__device__ Camera()
{
lowerLeftCorner = Vec3(-2.0, -1.0, -1.0);
horizontal = Vec3( 4.0, 0.0, 0.0);
vertical = Vec3( 0.0, 2.0, 0.0);
origin = Vec3( 0.0, 0.0, 0.0);
}
__device__ Ray GetRay(double u, double v) const
{
Vec3 direction = lowerLeftCorner + u * horizontal + v * vertical - origin;
return Ray(origin, direction);
}
private:
Vec3 origin;
Vec3 lowerLeftCorner;
Vec3 horizontal;
Vec3 vertical;
};
C++
복사
3.3 CPU 대비 제거된 기능
•
FOV / lookAt 설정
•
피사계 심도(Defocus blur)
•
픽셀 샘플 오프셋(카메라 내부가 아니라 커널에서 cuRAND로 처리)
4. kernel.cu: 렌더링 파이프라인(전체 흐름)
4.1 실행 순서
1.
createWorld<<<1,1>>> : GPU에서 Sphere 2개 + HittableList + Camera 생성
2.
renderInit<<<blocks, threads>>> : 픽셀당 curandState 초기화
3.
render<<<blocks, threads>>> : 픽셀당 N회 샘플링 후 평균
4.
CPU에서 프레임버퍼를 읽어 PPM 저장
5.
freeWorld<<<1,1>>> : GPU 오브젝트 해제
4.2 렌더 커널 핵심(멀티샘플 평균)
__global__ void render(
Color* frameBuffer,
int maxX,
int maxY,
int numSamples,
Camera** camera,
Hittable** world,
curandState* randState)
{
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 + i;
curandState local = randState[pixelIndex];
Color col(0.0, 0.0, 0.0);
for (int s = 0; s < numSamples; ++s) {
// 픽셀 내부 랜덤 위치로 지터링(jittering)
double u = (double(i) + double(curand_uniform(&local))) / double(maxX);
double v = (double(j) + double(curand_uniform(&local))) / double(maxY);
Ray r = (*camera)->GetRay(u, v);
col += RayColor(r, world);
}
randState[pixelIndex] = local;
frameBuffer[pixelIndex] = col / double(numSamples);
}
C++
복사
설정 | 값 | 설명 |
numSamples | 100 | 픽셀당 샘플 수 |
랜덤 오프셋 | curand_uniform | (0, 1] 범위 난수로 픽셀 내부 위치 샘플링 |
UV 계산 | (i + random) / maxX | 픽셀 내부 지터링된 좌표 |
5. GPU 메모리 사용량(대략)
변수 | 타입 | 크기(대략) | 용도 |
frameBuffer | Color* | 1440×720 × sizeof(Color) | 픽셀 색상 저장 |
randState | curandState* | 1440×720 × sizeof(curandState) | 픽셀당 난수 상태 |
list | Hittable** | 2 pointers | 구체 포인터 배열 |
world | Hittable** | 1 pointer | HittableList 포인터 |
camera | Camera** | 1 pointer | Camera 포인터 |
렌더링 결과
•
해상도: 1440 × 720
•
샘플 수: 100 samples/pixel
•
오브젝트: 구체 2개 (중앙 r=0.5, 바닥 r=100)
•
효과: 구체 가장자리가 부드럽게 블렌딩(안티 앨리어싱)
•
출력 파일: output.ppm