목차: Series 1 – Index + Methodology (Link)
이전 글: Series 4 – Classing Up the GPU (Link)
다음 글: Series 6 – It’s All Random (Link)
이번 장에서는 여러 개의 Sphere(구) 정보를 GPU 메모리에 저장하고 해당 정보를 활용해서 Sphere를 Rendering 할 예정이다. 이러한 과정에서 GPU 메모리를 관리하는 방법에 관해서 설명한다. 먼저 기존에 CPU에서 작성한 hitable, hitablelist, sphere Class를 GPU가 사용할 수 있도록 변경해야 한다. 아래 코드 1, 2, 3은 기존 CPU 코드에서 GPU가 호출할 수 있도록 함수 이름 앞에 “__device__”를 추가하였다.
코드 1: hitable.h 코드
#ifndef MKHITABLE_H
#define MKHITABLE_H
#include "mkVec3.h"
//MK: Ray와 Hit한 Object의 위치를 파악하기 위해 사용
struct hitRecord{
float t;
vec3 p;
vec3 normal;
};
//MK: (코드 1-1) GPU에서 호출 가능하도록 __device__를 추가함
class hitable{
public:
__device__ virtual bool hit(const ray &r, float tMin, float tMax, hitRecord &rec) const = 0;
};
#endif
코드 2: Hitablelist.h 코드
#ifndef HITABLELIST_H
#define HITABLELIST_H
#include "mkHitable.h"
//MK: 모든 Hitable Object를 리스트로 가지고 있음
//MK: Ray와 모든 Ojbect의 Hit(Intersection)여부를 판단함
//MK: (코드 2-1) GPU가 호출가능하도록 함수 이름 앞에 __device__를 추가함
class hitableList: public hitable{
public:
__device__ hitableList(){}
__device__ hitableList(hitable **l, int n){
list = l;
listSize = n;
}
__device__ virtual bool hit(const ray &r, float tMin, float tMax, hitRecord &rec) const {
hitRecord tempRec;
bool hitAnything = false;
double closestSoFar = tMax;
for(int i = 0; i < listSize; i++){
if(list[i]->hit(r, tMin, closestSoFar, tempRec)){
hitAnything = true;
closestSoFar = tempRec.t;
rec = tempRec;
}
}
return hitAnything;
}
private:
hitable **list;
int listSize;
};
#endif코드 3: sphere.h 코드
#ifndef MKSPHERE_H
#define MKSPHERE_H
#include "mkHitable.h"
//MK: Sphere Object를 생성하여 Ray와 Hit여부를 판단하는 부분
//MK: (코드 3-1) GPU에서 호출 가능하도록 __device__를 추가함
class sphere: public hitable{
public:
__device__ sphere(){}
__device__ sphere(vec3 cen, float r) : center(cen), radius(r) {}
__device__ virtual bool hit(const ray &r, float tMin, float tMax, hitRecord &rec) const {
vec3 oc = r.origin() - center;
float a = dot(r.direction(), r.direction());
float b = dot(oc, r.direction());
float c = dot(oc, oc) - radius * radius;
float discriminant = b*b - a*c;
if(discriminant > 0){
float temp = (-b - sqrt(discriminant))/a;
//MK: 구의 가까운 부분 부터 Hit여부를 판단함
if(temp < tMax && temp > tMin){
rec.t = temp;
rec.p = r.pointAtParameter(rec.t);
rec.normal = (rec.p - center) / radius;
return true;
}
temp = (-b + sqrt(discriminant))/a;
if(temp < tMax && temp > tMin){
rec.t = temp;
rec.p = r.pointAtParameter(rec.t);
rec.normal = (rec.p - center) / radius;
return true;
}
}
return false;
}
private:
vec3 center;
float radius;
};
#endif이제 Sphere를 GPU 메모리 공간에 할당하는 코드를 작성할 차례이다. 아래 코드 4는 Main 파일 코드이다. 코드 4-1에 dList, dWorld라는 Array를 cudaMalloc 함수를 사용해서 GPU 메모리에 할당한다. dList의 경우 hitable 포인터를 2개 저장할 수 있는 공간을 할당하고, dWorld의 경우 hitable 포인터 하나를 저장할 수 있는 공간을 할당한다. cudaMalloc의 경우 GPU에서만 Access 가능한 메모리 공간을 할당한다. cudaMalloc으로 할당된 메모리 공간은 CPU에서 접근이 불가능하다. 만약 GPU 메모리 공간에 저장된 값을 읽어야 하는 경우 CPU 메모리(Main Memory)로 값을 먼저 복사하는 작업을 수행해야 한다.
코드 4: Main 파일 코드
#include <fstream>
#include "mkRay.h"
#include <time.h>
#include "mkSphere.h"
#include "mkHitablelist.h"
#include <float.h>
using namespace std;
//MK: FB 사이즈
int nx = 1200;
int ny = 600;
// limited version of checkCudaErrors from helper_cuda.h in CUDA examples
//MK: #val은 val 전체를 String으로 Return 함 (출처 3)
#define checkCudaErrors(val) check_cuda( (val), #val, __FILE__, __LINE__ )
//MK: Error 위치를 파악하기 위해서 사용
void check_cuda(cudaError_t result, char const *const func, const char *const file, int const line) {
if (result) {
cerr << "MK: CUDA ERROR = " << static_cast<unsigned int>(result) << " at " << file << ":" << line << " '" << func << "' \n";
// Make sure we call CUDA Device Reset before exiting
cudaDeviceReset();
exit(99);
}
}
//MK: (코드 4-2) Sphere 2개를 추가하는 코드
//MK: 1개의 Thread만 연산을 수행할 수 있도록 if문을 추가함
__global__ void mkCreateWorld(hitable **dList, hitable **dWorld){
if(threadIdx.x == 0 && blockIdx.x == 0){
*(dList) = new sphere(vec3(0, 0, -1), 0.5);
*(dList + 1) = new sphere(vec3(0, -100.5, -1), 100);
*dWorld = new hitableList(dList, 2);
}
}
//MK: (코드 4-7) 여러개의 Sphere Hit 여부를 판단하여 색상을 결정하도록 코드 변경
__device__ vec3 color(const ray &r, hitable **dWorld){
hitRecord rec;
vec3 ret = vec3(0, 0, 0);
if((*dWorld)->hit(r, 0.0, FLT_MAX, rec)){
ret = 0.5f * vec3(rec.normal.x() + 1.0f, rec.normal.y() + 1.0f, rec.normal.z() + 1.0f);
}
else{
vec3 unitDirection = unitVector(r.direction());
float t = 0.5f * (unitDirection.y() + 1.0f);
ret = (1.0 - t) * vec3(1.0, 1.0, 1.0) + t * vec3(0.5, 0.7, 1.0);
}
return ret;
}
//MK: (코드 4-6) 여러개의 Sphere의 색상을 결정하기 위해서 코드 변경
__global__ void mkRender(vec3 *fb, int max_x, int max_y, vec3 lowerLeftCorner, vec3 horizontal, vec3 vertical, vec3 origin, hitable **dWorld) {
//MK: Pixel 위치 계산을 위해 ThreadId, BlockId를 사용함
int i = threadIdx.x + blockIdx.x * blockDim.x;
int j = threadIdx.y + blockIdx.y * blockDim.y;
//MK: 계산된 Pixel 위치가 FB사이즈 보다 크면 연산을 수행하지 않음
if((i >= max_x) || (j >= max_y)){
return;
}
//MK: FB Pixel 값 계산
int pixel_index = j*max_x + i;
float u = float(i)/float(max_x);
float v = float(j)/float(max_y);
ray r(origin, lowerLeftCorner + u*horizontal + v*vertical);
fb[pixel_index] = color(r, dWorld);
}
//MK: (코드 4-4) mkCreateWorld에서 생성한 클래스 제거
__global__ void mkFreeWorld(hitable **dList, hitable **dWorld){
if(threadIdx.x == 0 && blockIdx.x == 0){
delete *(dList);
delete *(dList + 1);
delete *dWorld;
}
}
int main() {
//MK: Thread Block 사이즈
int tx = 8;
int ty = 8;
cout << "MK: Rendering a " << nx << "x" << ny << " Image ";
cout << "MK: in " << tx << "x" << ty << " Thread Blocks.\n";
clock_t start, stop;
start = clock();
int num_pixels = nx*ny;
size_t fb_size = 3*num_pixels*sizeof(float);
//MK: FB 메모리 할당 (cudaMallocManaged 는 Unitifed Memory를 사용 할 수 있도록 함)
//MK: 필요에 따라 CPU/GPU에서 GPU/CPU로 데이터를 복사함
vec3 *fb;
checkCudaErrors(cudaMallocManaged((void **)&fb, fb_size));
//MK: (코드 4-1) Sphere을 여러개 추가하기 위해서 메모리 할당을 진행
hitable **dList;
hitable **dWorld;
checkCudaErrors(cudaMalloc((void **) &dList, 2 * sizeof(hitable *)));
checkCudaErrors(cudaMalloc((void **) &dWorld, sizeof(hitable *)));
//MK: (코드 4-3) Sphere를 생성하는 함수를 호출함
mkCreateWorld<<<1, 1>>>(dList, dWorld);
checkCudaErrors(cudaGetLastError());
checkCudaErrors(cudaDeviceSynchronize());
//MK: GPU (CUDA) 연산을 위해서 Thread Block, Grid 사이즈 결정
dim3 blocks(nx/tx+1,ny/ty+1);
dim3 threads(tx,ty);
//MK: CUDA 함수 호출
mkRender<<<blocks, threads>>>(fb, nx, ny,
vec3(-2.0, -1.0, -1.0),
vec3(4.0, 0.0, 0.0),
vec3(0.0, 2.0, 0.0),
vec3(0.0, 0.0, 0.0),
dWorld);
checkCudaErrors(cudaGetLastError());
//MK: CUDA 연산이 완료되길 기다림
checkCudaErrors(cudaDeviceSynchronize());
//MK: 연산 시간과 끝 부분을 계산하여서 연산 시간을 측정함
//MK: CPU 코드와 동일하게 결과를 파일에 작성
string fileName = "Ch5_gpu.ppm";
ofstream writeFile(fileName.data());
if(writeFile.is_open()){
writeFile.flush();
writeFile << "P3\n" << nx << " " << ny << "\n255\n";
for (int j = ny-1; j >= 0; j--) {
for (int i = 0; i < nx; i++) {
size_t pixel_index = j*nx + i;
int ir = int(255.99 * fb[pixel_index].r());
int ig = int(255.99 * fb[pixel_index].g());
int ib = int(255.99 * fb[pixel_index].b());
writeFile << ir << " " << ig << " " << ib << "\n";
}
}
writeFile.close();
}
//MK: (코드 4-5) 사용한 메모리를 제거함
mkFreeWorld<<<1, 1>>>(dList, dWorld);
checkCudaErrors(cudaGetLastError());
checkCudaErrors(cudaFree(dList));
checkCudaErrors(cudaFree(dWorld));
checkCudaErrors(cudaFree(fb));
//MK: 연산 시간과 끝 부분을 계산하여서 연산 시간을 측정함
stop = clock();
double timer_seconds = ((double)(stop - start)) / CLOCKS_PER_SEC;
cout << "MK: GPU (CUDA) Took " << timer_seconds << " Seconds.\n";
return 0;
}코드 4-2는 Sphere Class를 생성하여 dList, dWorld Array에 추가하는 코드이다. 중간에 Class를 생성하여서 Array에 저장하기 이전에 threadIdx, blockIdx를 사용해서 1개의 Thread만 동작하도록 코드를 작성하였다. Sphere 정보는 오직 한 번만 생성하면 된다. 한번 생성된 Sphere를 여러 개의 Thread가 사용하면 되므로 IF 문을 사용해서 하나의 Thread만 연산을 수행하도록 코드를 작성한다.
코드 4-3은 mkCreateWorld 함수를 호출하는 코드이다. 1개의 Thread만 dList, dWorld를 생성하면 되므로 BlockDim, GridDim을 모두 1로 설정하여서 creatWorld 함수를 호출한다.
해당 Sphere Class를 사용해서 새로운 이미지를 Rendering 하기 이전에 사용한 메모리를 Free 하는 부분부터 설명할 예정이다. 코드 4-4는 mkCreateWorld에서 생성한 Class를 제거하는 코드이다. 코드 4-2와 동일하게 threadIdx, blockIdx를 사용해서 1개의 Thread만 해당 코드를 수행하도록 코드를 작성하였다. 코드 4-5 mkFreeWorld를 호출하고 cudaFree를 사용해서 GPU 메모리를 Free 하는 코드이다.
이제 생성한 Sphere Class 정보를 사용해서 색상을 결정하는 코드이다. 코드 4-6은 mkRender 함수 코드를 수정한 것이다. 기존 코드와 큰 차이가 없다. 코드 4-7은 Sphere와 Ray의 Hit 여부를 판단해서 색상을 결정하는 코드이다.
위 그림 1은 앞서 작성한 코드를 컴파일해서 실행하면 나타나는 결과 이미지이다. 코드를 원하시는 분이 있을 수도 있어서 아래 출처 3에 CPU/GPU 코드를 공유하였다.
성능 측정 결과
그림 2는 CPU와 GPU의 연산 시간을 정리한 결과이다. CPU는 대략 0.32초 정도의 계산 시간이 필요하다. GPU의 경우 대략 0.19초 정도의 연산 시간이 필요하다. GPU가 대략 1.7배 정도 빠른 성능을 가진다.
출처
- https://devblogs.nvidia.com/accelerated-ray-tracing-cuda/
- http://www.realtimerendering.com/raytracing/Ray%20Tracing%20in%20a%20Weekend.pdf
- https://github.com/mkblog-cokr/CudaRTSeries