[CUDA RT Series 6] It’s All Random

by

목차: Series 1 – Index + Methodology (Link)
이전 글: Series 5 – Manage Your Memory (Link)
다음 글: Series 7 – Iteration vs. Recursion (Link)

이번 글에서는 Anti-Aliasing(AA)를 구현한다. AA를 구현하기 위해서 Random 함수를 사용한다. CUDA는 cuRAND 라이브러리를 사용해서 Random 수를 생성한다. cuRAND는 CUDA Random Number Generation 라이브러리이다. cuRAND 라이브러리를 사용하여 AA를 구현하기 이전에 먼저 Camera Class를 먼저 생성한다. 아래 코드 1은 기존 CPU 코드를 복사하여서 GPU에서 실행할 수 있도록 __device__를 추가한 코드이다.

MK: AA를 구현하기 위해서 Camera Class를 생성한 것은 아니다. 출처 1의 글은 CPU 코드로 작성된 코드를 GPU(CUDA) 코드로 작성하는 방법에 관해서 설명한다. CPU에서 AA 구현 이전에 Camera에 대한 부분을 먼저 설명한다. Camera 코드는 앞에서 설명했던 부분과 같이 CPU 코드에 __device__ 를 추가하면 되는 관계로 따로 설명하는 글을 작성하지 않을 것으로 판단된다.

코드 1: Camera Class 코드

#ifndef MKCAMERA_H
#define MKCAMERA_H

#include "mkRay.h"

//MK: (코드 1) CPU Camera 코드를 복사후 GPU에서 실행하기 위해서 함수 이름 앞에 __device__를 추가함
class camera {
	public:
		__device__ camera() {
			lower_left_corner = 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 get_ray(float u, float v) { 
			return ray(origin, lower_left_corner + u*horizontal + v*vertical - origin); 
		}
	
	private:
        vec3 origin;
        vec3 lower_left_corner;
        vec3 horizontal;
        vec3 vertical;
};

#endif

이제 cuRAND를 사용해서 AA를 구현할 차례이다. 아래 코드 2는 cuRAND를 사용해서 AA를 구현한 Main 함수 코드이다. 

코드 2: cuRAND 라이브러리를 사용한 AA 구현 코드

#include <fstream>
#include "mkRay.h"
#include <time.h>
#include "mkSphere.h"
#include "mkHitablelist.h"
#include <float.h>
//MK: (코드 2-1) Camera와 curand를 사용하기 위해서 추가함
#include <curand_kernel.h>
#include "mkCamera.h"

using namespace std;

//MK: FB 사이즈
//MK: (코드 2-2) Sample 수(ns)를 추가함
int nx = 1200;
int ny = 600;
int ns = 100;

// limited version of checkCudaErrors from helper_cuda.h in CUDA examples
#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: (코드 2-3) 새로추가한 Camera를 생성하기 위해서 코드를 수정함
__global__ void mkCreateWorld(hitable **dList, hitable **dWorld, camera **dCamera){
    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);
    	*dCamera = new camera();
    }
}

__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: (코드 2-4) curand_uniform을 사용해서 AA(Anti-Aliasing)을 구현함
__global__ void mkRender(vec3 *fb, int max_x, int max_y, int num_sample, camera **cam, 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;
 	
    curandState rand_state;
    //MK: (코드 2-5) 원래 코드는 아래 부분을 사용한다. 하지만 메모리 관련 에러가 발생하여서 약간 수정하였다.
    //curand_init(1984, pixel_index, 0, &rand_state);
    curand_init(pixel_index, 0, 0, &rand_state);
    vec3 col(0, 0, 0);
    for(int s = 0; s < num_sample; s++){
	float u = float(i + curand_uniform(&rand_state))/float(max_x);
   	float v = float(j + curand_uniform(&rand_state))/float(max_y);
	ray r = (*cam)->get_ray(u, v);
	col += color(r, dWorld);
    }
    fb[pixel_index] = col/float(num_sample);
}

//MK: (코드 2-6) 새로 생성한 Camera를 제거하기 위해서 코드를 수정함
__global__ void mkFreeWorld(hitable **dList, hitable **dWorld, camera **dCamera){
    if(threadIdx.x == 0 && blockIdx.x == 0){
        delete *(dList);
	delete *(dList + 1);
	delete *dWorld;
	delete *dCamera;
    }
}

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: (코드 2-7) 새로운 Camera를 사용하기 위해서 GPU 메모리 공간을 할당함
    hitable **dList;
    hitable **dWorld;
    camera **dCamera;
    checkCudaErrors(cudaMalloc((void **) &dList, 2 * sizeof(hitable *)));
    checkCudaErrors(cudaMalloc((void **) &dWorld, sizeof(hitable *)));
    checkCudaErrors(cudaMalloc((void **) &dCamera, sizeof(camera *)));
		
    //MK: (코드 2-8) Camera를 생성하기 위해서 코드를 수정함
    mkCreateWorld<<<1, 1>>>(dList, dWorld, dCamera);
    checkCudaErrors(cudaGetLastError());
    checkCudaErrors(cudaDeviceSynchronize());

    //MK: GPU (CUDA) 연산을 위해서 Thread Block, Grid 사이즈 결정
    dim3 blocks(nx/tx+1,ny/ty+1);
    dim3 threads(tx,ty);
    //MK: CUDA 함수 호출
    //MK: (코드 2-9) Camera를 사용해서 Rendering하기 위해서 코드를 수정함
    mkRender<<<blocks, threads>>>(fb, nx, ny, ns, dCamera, dWorld);
    checkCudaErrors(cudaGetLastError());
    //MK: CUDA 연산이 완료되길 기다림
    checkCudaErrors(cudaDeviceSynchronize());
    //MK: 연산 시간과 끝 부분을 계산하여서 연산 시간을 측정함 

    //MK: CPU 코드와 동일하게 결과를 파일에 작성
    string fileName = "Ch6_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: (코드 2-10) 모두 사용한 Camera Class를 제거하고 메모리를 Release하기 위해서 코드를 수정함
    mkFreeWorld<<<1, 1>>>(dList, dWorld, dCamera);
    checkCudaErrors(cudaGetLastError());
    checkCudaErrors(cudaFree(dList));
    checkCudaErrors(cudaFree(dWorld));
    checkCudaErrors(cudaFree(fb));
    checkCudaErrors(cudaFree(dCamera));
	
    //MK: 연산 시간과 끝 부분을 계산하여서 연산 시간을 측정함 
    stop = clock();
    double timer_seconds = ((double)(stop - start)) / CLOCKS_PER_SEC;
    cout << "MK: GPU (CUDA) Took " << timer_seconds << " Seconds.\n";

    return 0;
}

먼저 새로 생성한 Camera Class를 사용하기 위해서 GPU 메모리를 할당하고 Camera Class를 생성하고 제거하는 부분부터 코드를 작성한다. 코드 2-7, 2-3, 2-6, 2-10은 GPU 메모리를 할당하고, Camera Class를 생성하고, 생성한 Class를 제거하고, GPU 메모리를 제거(Free/Release)하는 부분의 코드이다. 이전 글에서 설명한 내용인 관계로 상세한 설명은 제외하였다.

cuRAND 라이브러리를 사용하기 위해서 “curand_kernel.h” 해더를 추가한다 (코드 2-1). 다음은 Pixel 당 Sampling을 진행할 횟수를 지정한다. 저의 경우 100회로 지정하였다 (코드 2-2). 이제 cuRAND 라이브러리를 사용해서 Random 숫자를 생성하여 Pixel 값을 계산하는 코드를 작성한다. 코드 2-4는 curand_init(..) 함수와 curand_uniform(…)을 사용해서 Random 수를 생성하는 코드이다. 생성된 Random 수를 사용해서 하나의 Pixel 값을 총 100번 계산하다. 100번의 값을 평균하여서 최종 Pixel 값을 결정한다. AA 구현 설명은 이전에 작성한 글(출처 3)을 참조하면 된다. 

MK: 출처 1의 코드와 약간의 차이가 있다. 출처 1에서는 curand_init(…)을 다른 함수에서 실행한다. 출처 1에 따르면 curand_init(…) 연산 시간을 총 실행 시간에서 제외하기 위해서라고 한다. 해당 시간 차이가 크지 않을 것으로 판단하여 함수를 제거하였다 (코드 2-4 참조).

위에 작성한 코드를 컴파일(NVCC)하여 실행하면 아래 그림 1 결과를 확인 할 수 있다. 지난 글에서 작성한 코드의 결과와 거의 동일하다. 하지만, 구의 끝부분(Edge)을 보면 이전 결과보다 부드러운 형태로 변경된 것을 확인 할 수 있다.

그림 1: 결과 이미지

성능 측정 결과

그림 2: CPU vs. GPU 성능 결과

그림 2는 CPU/GPU 성능을 비교한 결과이다. 이전까지는 GPU가 대략 2~3배 수준으로 빠른 성능을 보였다. 하지만, AA의 경우 계산량이 급격히 증가하게 된다. AA에서 모든 Pixel 값을 계산하기 위해서 총 1200 x 600 x 100번은 Pixel 계산이 필요하다. CPU의 경우 하나의 Thread가 모든 Pixel 값 연산을 순차적으로 수행한다. 반면, GPU의 경우 1200 x 600개의 Thread가 100번을 연산을 수행한다. 그 결과 GPU가 대략 60배 정도 빨리 연산을 완료하는 것을 확인 할 수 있다.

출처

  1. https://devblogs.nvidia.com/accelerated-ray-tracing-cuda/
  2. http://www.realtimerendering.com/raytracing/Ray%20Tracing%20in%20a%20Weekend.pdf
  3. https://mkblog.co.kr/2019/06/06/rt-in-one-weekend-series-7-anti-aliasing/
  4. https://github.com/mkblog-cokr/CudaRTSeries

Leave a Comment