Bounding Volume Hierarchies (경계 볼륨 계층 구조)

if (레이가 경계 볼륨에 맞음) return 내부 물체들에 맞는지 여부 else return false

if (보라색에 맞음) hit0 = 파란색 안의 물체들에 맞는지 hit1 = 빨간색 안의 물체들에 맞는지 if (hit0 or hit1) return true + 더 가까운 히트 정보 return false

P(t) = Q + t·d

interval_x ← compute_intersection_x(ray, x0, x1) interval_y ← compute_intersection_y(ray, y0, y1) return overlaps(interval_x, interval_y)

interval_x ← compute_intersection_x(ray, x0, x1) interval_y ← compute_intersection_y(ray, y0, y1) interval_z ← compute_intersection_z(ray, z0, z1) return overlaps(interval_x, interval_y, interval_z)

bool overlaps(t1, t2) t_min ← max(t1.min, t2.min) t_max ← min(t1.max, t2.max) return t_min < t_max

__host__ __device__ double Clamp(double value) const { if (value < Min) return Min; if (value > Max) return Max; return value; } // 구간을 delta만큼 양쪽으로 넓힌다. // 두께가 0인 AABB(축에 완전히 평행한 면)에서 생기는 grazing/NaN 문제를 막는 패딩용. __host__ __device__ Interval Expand(double delta) const { double padding = delta / 2.0; return Interval(Min - padding, Max + padding); }

#pragma once #ifndef AABB_H #define AABB_H #include"Ray.h" #include"Interval.h" // 3차원 AABB = x/y/z 세 슬랩(구간)의 교집합. // 레이가 이 상자를 통과하는지 빠르게 판정하는 것이 BVH 가속의 핵심이다. class Aabb { public: Interval X; Interval Y; Interval Z; // 기본 AABB는 비어 있다(Interval 기본 생성자가 빈 구간을 만든다). __host__ __device__ Aabb() {} __host__ __device__ Aabb(const Interval& x, const Interval& y, const Interval& z) : X(x), Y(y), Z(z) { } // 두 점 a, b를 상자의 양 극점으로 보고 만든다(좌표 대소 순서 신경 안 써도 됨). __host__ __device__ Aabb(const Point3& a, const Point3& b) { X = (a[0] <= b[0]) ? Interval(a[0], b[0]) : Interval(b[0], a[0]); Y = (a[1] <= b[1]) ? Interval(a[1], b[1]) : Interval(b[1], a[1]); Z = (a[2] <= b[2]) ? Interval(a[2], b[2]) : Interval(b[2], a[2]); } // (Listing 13) 두 AABB를 모두 감싸는 AABB. __host__ __device__ Aabb(const Aabb& box0, const Aabb& box1) : X(box0.X, box1.X) , Y(box0.Y, box1.Y) , Z(box0.Z, box1.Z) { } __host__ __device__ const Interval& AxisInterval(int n) const { if (n == 1) return Y; if (n == 2) return Z; return X; } // 슬랩 방법으로 레이-상자 교차 여부만 판정한다(히트 지점/법선 불필요). // Ray 접근자가 __device__ 전용이라 이 함수도 __device__ 전용이다. // ※ 위 콜아웃 참고: 루프/조기 return 대신 분기 없는 fmin/fmax 형태. __device__ bool Hit(const Ray& ray, Interval rayT) const { const Point3& o = ray.Origin(); const Vector3& d = ray.Direction(); double tMin = rayT.Min; double tMax = rayT.Max; // X 축 슬랩 double invDx = 1.0 / d[0]; double tx0 = (X.Min - o[0]) * invDx; double tx1 = (X.Max - o[0]) * invDx; tMin = fmax(tMin, fmin(tx0, tx1)); tMax = fmin(tMax, fmax(tx0, tx1)); // Y 축 슬랩 double invDy = 1.0 / d[1]; double ty0 = (Y.Min - o[1]) * invDy; double ty1 = (Y.Max - o[1]) * invDy; tMin = fmax(tMin, fmin(ty0, ty1)); tMax = fmin(tMax, fmax(ty0, ty1)); // Z 축 슬랩 double invDz = 1.0 / d[2]; double tz0 = (Z.Min - o[2]) * invDz; double tz1 = (Z.Max - o[2]) * invDz; tMin = fmax(tMin, fmin(tz0, tz1)); tMax = fmin(tMax, fmax(tz0, tz1)); return tMax > tMin; } // (Listing 21) 가장 긴 축의 인덱스 — 분할 축 선택용. __host__ __device__ int LongestAxis() const { if (X.Size() > Y.Size()) return X.Size() > Z.Size() ? 0 : 2; else return Y.Size() > Z.Size() ? 1 : 2; } }; #endif

#include"Ray.h" #include"AABB.h" class Material; // struct HitRecord { ... } // 기존과 동일 class Hittable { public: __device__ virtual bool Hit( const Ray& ray, double tMin, double tMax, HitRecord& hitRecord) const = 0; // 이 객체를 감싸는 AABB. 움직이는 물체는 운동 전체 구간(time 0~1)을 감싼다. __device__ virtual Aabb BoundingBox() const = 0; // BVH 반복 순회에서 내부 노드/잎 구분용(기본 false = 잎/일반 오브젝트). __device__ virtual bool IsBvhNode() const { return false; } };

__device__ Sphere(const Point3& center, double radius, Material* material) : mCenter(center) , mRadius(radius) , mMaterial(material) { // 반지름 벡터로 중심 ± r 두 극점을 잡아 경계 상자를 만든다. Vector3 rvec(radius, radius, radius); mBBox = Aabb(center - rvec, center + rvec); } // ... Hit() 은 기존과 동일 ... __device__ Aabb BoundingBox() const override { return mBBox; } private: Point3 mCenter; double mRadius; Material* mMaterial; Aabb mBBox; // 추가

__device__ MovingSphere( Point3 center0, Point3 center1, double time0, double time1, double radius, Material* material) : mCenter0(center0) , mCenter1(center1) , mTime0(time0) , mTime1(time1) , mRadius(radius) , mMaterial(material) { // 운동 구간 전체를 감싸려면 time0 위치 박스와 time1 위치 박스를 모두 포함해야 한다. Vector3 rvec(radius, radius, radius); Aabb box0(center0 - rvec, center0 + rvec); Aabb box1(center1 - rvec, center1 + rvec); mBBox = Aabb(box0, box1); } // ... Hit() 은 기존과 동일 ... __device__ Aabb BoundingBox() const override { return mBBox; } private: // ... 기존 멤버 ... Aabb mBBox; // 운동 전체를 감싸는 박스

// 두 구간을 빈틈없이 감싸는 구간을 생성한다(AABB 합집합용). __host__ __device__ Interval(const Interval& a, const Interval& b) : Min(a.Min <= b.Min ? a.Min : b.Min) , Max(a.Max >= b.Max ? a.Max : b.Max) { }

// AABB.h 안 — 두 AABB를 모두 감싸는 AABB __host__ __device__ Aabb(const Aabb& box0, const Aabb& box1) : X(box0.X, box1.X) , Y(box0.Y, box1.Y) , Z(box0.Z, box1.Z) { }

__device__ HittableList(Hittable** list, int count) : mList(list), mCount(count) { // 자식들의 경계 상자를 합쳐 리스트 전체의 경계 상자를 구한다. for (int i = 0; i < count; i++) mBBox = Aabb(mBBox, list[i]->BoundingBox()); } // ... Hit() 은 기존과 동일 ... __device__ Aabb BoundingBox() const override { return mBBox; } private: Hittable** mList; int mCount; Aabb mBBox; // 추가

#pragma once #ifndef BVH_NODE_H #define BVH_NODE_H #include"Hittable.h" #include"AABB.h" class BvhNode : public Hittable { public: __device__ BvhNode() {} // objects : 잎 포인터 배열. [start,end) 구간을 제자리 정렬한다. // nodeStorage : 생성된 모든 BvhNode 등록(해제용) — shared_ptr 대체. // nodeCount : nodeStorage 등록 개수(공유 카운터). __device__ BvhNode( Hittable** objects, int start, int end, Hittable** nodeStorage, int* nodeCount) { // (Listing 20) span의 bbox를 먼저 만들고, 가장 긴 축으로 분할. mBBox = Aabb(); for (int i = start; i < end; i++) mBBox = Aabb(mBBox, objects[i]->BoundingBox()); int axis = mBBox.LongestAxis(); int objectSpan = end - start; if (objectSpan == 1) { // 잎이 하나면 양쪽 자식에 같은 잎(널 포인터 검사 회피). mLeft = mRight = objects[start]; } else if (objectSpan == 2) { mLeft = objects[start]; mRight = objects[start + 1]; } else { DeviceSort(objects, start, end, axis); // std::sort 대체 int mid = start + objectSpan / 2; BvhNode* leftNode = new BvhNode(objects, start, mid, nodeStorage, nodeCount); BvhNode* rightNode = new BvhNode(objects, mid, end, nodeStorage, nodeCount); nodeStorage[(*nodeCount)++] = leftNode; // 해제 레지스트리 등록 nodeStorage[(*nodeCount)++] = rightNode; mLeft = leftNode; mRight = rightNode; } } // (Listing 15) hit — 재귀 대신 고정 크기 명시적 스택으로 반복 순회. __device__ bool Hit( const Ray& ray, double tMin, double tMax, HitRecord& hitRecord) const override { const BvhNode* stack[32]; // 균형 분할이라 트리 높이 ~log2(N), 32면 충분 int sp = 0; const BvhNode* node = this; bool hitAnything = false; double closest = tMax; while (true) { // 현재 노드의 박스에 맞을 때만 자식을 살펴본다. if (node->mBBox.Hit(ray, Interval(tMin, closest))) { const BvhNode* next = nullptr; Hittable* kids[2] = { node->mLeft, node->mRight }; for (int c = 0; c < 2; c++) { Hittable* kid = kids[c]; if (kid->IsBvhNode()) { // 내부 노드: 하나는 바로 내려가고 나머지는 스택에 보관. const BvhNode* bn = static_cast<const BvhNode*>(kid); if (next == nullptr) next = bn; else if (sp < 32) stack[sp++] = bn; } else { // 잎(primitive): 직접 검사. 더 가까우면 closest 갱신. if (kid->Hit(ray, tMin, closest, hitRecord)) { hitAnything = true; closest = hitRecord.T; } } } if (next != nullptr) { node = next; continue; } } if (sp == 0) break; node = stack[--sp]; } return hitAnything; } __device__ Aabb BoundingBox() const override { return mBBox; } __device__ bool IsBvhNode() const override { return true; } private: Hittable* mLeft; Hittable* mRight; Aabb mBBox; // (Listing 18) 분할 축의 bbox 최솟값으로 a < b 비교. __device__ static bool BoxCompare(const Hittable* a, const Hittable* b, int axisIndex) { Interval aAxis = a->BoundingBox().AxisInterval(axisIndex); Interval bAxis = b->BoundingBox().AxisInterval(axisIndex); return aAxis.Min < bAxis.Min; } // std::sort 대체: [start,end) 구간을 분할 축 기준 제자리 삽입 정렬(단일 스레드). __device__ static void DeviceSort(Hittable** objects, int start, int end, int axis) { for (int i = start + 1; i < end; i++) { Hittable* key = objects[i]; int j = i - 1; while (j >= start && BoxCompare(key, objects[j], axis)) { objects[j + 1] = objects[j]; j--; } objects[j + 1] = key; } } }; #endif

// (재게시) 분할 축의 bbox 최솟값으로 a < b 비교. __device__ static bool BoxCompare(const Hittable* a, const Hittable* b, int axisIndex) { Interval aAxis = a->BoundingBox().AxisInterval(axisIndex); Interval bAxis = b->BoundingBox().AxisInterval(axisIndex); return aAxis.Min < bAxis.Min; }

// CreateWorld 시그니처: BVH 노드 배열/개수 출력 인자 추가 __global__ void CreateWorld( Hittable** list, Hittable** world, Camera** camera, int imageWidth, int imageHeight, curandState* randState, int* outCount, Hittable** bvhNodes, int* outNodeCount) { // ... 바닥/소형 구/대형 구를 list[0..i)에 배치 (기존과 동일) ... *outCount = i; // 실제 배치된 구 개수 // === BVH 빌드 === // list[0..i)를 BVH로 묶는다. 빌드 중 list는 분할 축 기준으로 제자리 정렬되지만, // 잎 포인터는 그대로 보존되므로 FreeWorld의 list[] 해제에는 영향이 없다. int nodeCount = 0; BvhNode* root = new BvhNode(list, 0, i, bvhNodes, &nodeCount); bvhNodes[nodeCount++] = root; // 루트도 해제 레지스트리에 등록 *outNodeCount = nodeCount; *world = root; // 기존 HittableList 대신 BVH 루트를 월드로 // ... 카메라 생성 (기존과 동일) ... }

__global__ void FreeWorld( Hittable** list, int numHittables, Hittable** bvhNodes, int numNodes, Hittable** world, Camera** camera) { if (threadIdx.x == 0 && blockIdx.x == 0) { for (int i = 0; i < numHittables; i++) delete list[i]; // 잎(primitive) for (int i = 0; i < numNodes; i++) delete bvhNodes[i]; // 내부 노드(*world == 루트도 여기 포함) delete *camera; } }

// BVH 순회는 재귀가 아니라 명시적 스택(BvhNode::Hit)을 쓰므로 추가 스택은 필요 없다. // (재귀로 두면 이 한도로도 일부 스레드에서 스택이 넘쳤다.) checkCudaErrors(cudaDeviceSetLimit(cudaLimitStackSize, 32768)); // ... // n개 잎에 대한 BVH 내부 노드 수는 최대 2n 미만 → 2*maxHittables 로 넉넉히 잡는다. Hittable** bvhNodes; checkCudaErrors(cudaMalloc((void**)&bvhNodes, 2 * maxHittables * sizeof(Hittable*))); int* d_numNodes; checkCudaErrors(cudaMallocManaged((void**)&d_numNodes, sizeof(int))); *d_numNodes = 0; CreateWorld<<<1, 1>>>(list, world, camera, imageWidth, imageHeight, randState2, d_numHittables, bvhNodes, d_numNodes); checkCudaErrors(cudaGetLastError()); checkCudaErrors(cudaDeviceSynchronize()); int numHittables = *d_numHittables; int numNodes = *d_numNodes; // ... RenderInit / Render ... // 해제 FreeWorld<<<1, 1>>>(list, numHittables, bvhNodes, numNodes, world, camera); checkCudaErrors(cudaGetLastError()); checkCudaErrors(cudaDeviceSynchronize()); checkCudaErrors(cudaFree(bvhNodes)); checkCudaErrors(cudaFree(d_numNodes)); // ... 나머지 cudaFree ...

__host__ __device__ int LongestAxis() const { // 바운딩 박스에서 가장 긴 축의 인덱스를 반환. if (X.Size() > Y.Size()) return X.Size() > Z.Size() ? 0 : 2; else return Y.Size() > Z.Size() ? 1 : 2; }