[CUDA RT Series 4] Classing Up the GPU (Single vs. Double Precision)

목차: Series 1 – Index + Methodology (Link)
이전 글: Series 3 – Adding Vectors (Link)
다음 글: Series 5 – Manage Your Memory (Link)

이번에는 GPU가 사용할 Class를 생성한다. GPU에서 연산할 Class는 기존 CPU 코드와 동일하게 작성한다. Series 3 (Adding Vecotrs)에서 작성한 것과 같이 GPU에서 호출할 함수 앞에 “__device__”를 추가하면 된다.  아래 코드 1은 GPU에서 호출할 수 있도록 작성한 Ray Class이다.

코드 1: Ray Class 코드

#ifndef MKRAY_H
#define MKRAY_H
 
#include "mkVec3.h"
 
//MK: (코드 1- 1) GPU에서 호출가능 하도록 클래스 내부 함수에 __device__를 추가함
class ray{
    public:
        __device__ ray(){}
        __device__ ray(const vec3 &a, const vec3 &b){
            A = a;
            B = b;
        }

        __device__ vec3 origin() const{
            return A;
        }

        __device__ vec3 direction() const{
            return B;
        }

        __device__ vec3 pointAtParameter(float t) const{
            return (A + t * B);
        }

    private:
        vec3 A;
        vec3 B;
};

#endif

다음은 Main 파일을 수정할 차례이다. 새로 생성한 Ray Class를 사용하도록 변경하고 하나의 Sphere(구)를 추가한다(코드 2-1, 코드 2-2 ). 아래 코드 2는 Ray Class를 사용하도록 Main 파일을 수정한 코드이다.

코드 2: Ray Class를 사용하기 위한 Main 코드

#include <fstream>
#include "mkRay.h"
#include <time.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: (코드 2-1) Ray와 Sphere(구)의 Hit여부를 판단
__device__ bool hitSphere(const vec3 &center, float radius, const ray &r){
    vec3 oc = r.origin() - center;
    float a = dot(r.direction(), r.direction());
    float b = 2.0f * dot(oc, r.direction());
    float c = dot(oc, oc) - radius * radius;
    float discriminant = b*b - 4.0f*a*c;
    return (discriminant >= 0);
}

//MK: (코드 2-2) 구의 Hit 여부에 따른 색상 결정
__device__ vec3 color(const ray &r){
    vec3 ret = vec3(1, 0, 0);
    if(hitSphere(vec3(0, 0, -1), 0.5, r)){
        return ret;
    }
    vec3 unitDirection = unitVector(r.direction());
    //MK: (코드 2-3) 중요 - Single/Double Precision 연산 설명
    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;
}

__global__ void mkRender(vec3 *fb, int max_x, int max_y, vec3 lowerLeftCorner, vec3 horizontal, vec3 vertical, vec3 origin) {
	//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);
}

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";

    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));

    clock_t start, stop;
    start = clock();
    //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));
    checkCudaErrors(cudaGetLastError());
	//MK: CUDA 연산이 완료되길 기다림
    checkCudaErrors(cudaDeviceSynchronize());
	//MK: 연산 시간과 끝 부분을 계산하여서 연산 시간을 측정함 

    //MK: CPU 코드와 동일하게 결과를 파일에 작성
	string fileName = "Ch4_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: 메모리 Free
    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;
}

Ray Class 내부의 함수와 동일하게 color와 hitSphere 함수 앞에 “__device__”를 추가하였다. 그 이외에 다른 부분은 CPU 코드와 거의 동일하다. 위 코드를 컴파일해서 실행하면 아래 그림 1 이미지 결과를 확인 할 수 있다.

그림 1: 결과 이미지

위 코드에서 가장 중요한 부분은 바로 아랫부분이다 (코드 2-3).

  • Single Precision: float t = 0.5f * (unitDirection.y() + 1.0f);
  • Double Precision: float t = 0.5 * (unitDirection.y() + 1.0);

위 코드 2개는 동일한 연산처럼 보인다. 하지만, 0.5f, 1.0f의 경우 모든 연산(곱하기, 더하기)이 Single Precision (32-Bit) 연산이다. 반면 아래 0.5, 1.0을 사용한 경우 모든 연산(곱하기, 더하기)이 Double Precision (64-Bit) 연산이다. 문제는 Double Precision 연산을 수행지만, 최종값은 Float (Single Precision) 결괏값이다. 다시 말해서 여기서 Double Precision 연산은 불필요한 연산인 셈이다.

만약 정교한 값을 사용해야 하는 경우 Double Precision 연산을 사용하면 된다. 보통 GPU의 경우 아주 많은 Single Precision (Float) 연산기를 가지고 있다. 반면, Double Precision 연산기 수는 상대적으로 아주 적은 수를 가지고 있다. 결과적으로 Double Precision 연산이 많은 Application의 경우 연산 시간이 상대적으로 늘어나게 되고 성능이 떨어진다. 최근에 나오는 NVIDIA GPU의 경우 Half Precision (16-Bit) 연산도 가능하다. NVIDIA 발표 자료에 따르면 Half Precision은 Single Precision의 2배 많은 연산을 동시에 수행할 수 있다고 한다 (출처 3).

실제로 코딩을 하다 보면 0.5f를 0.5로 작성하는 경우가 자주  발생한다. 그래서 “nvprof”라는 NVIDIA에서 제공하는 Profiling Tool을 사용해서 Single/Double Precision 명령어(Instruction) 수를 확인 할 수 있다. Linux에서 아래 명령어를 실행하면 Single/Double Precision 명령어 수를 확인 할 수 있다. “nvprof”은 CUDA Toolkit을 설치하면 같이 설치된다.

  • $ nvprof –metrics inst_fp_32,inst_fp_64 ./APP

MK: 제가 가지고 있는 컴퓨터에서는 위 CMD를 사용하면 에러가 발생한다. 인터넷에 검색하면 내가 가지고 있는 GPU는 위 CMD를 사용할 수 있는 것 같은데 정상적으로 동작을 하지 않는다. 다음에 문제를 해결하게 되면 결과 창을 정리해서 올려볼 계획이다. 우선은 Double Precision 연산이 발생하지 않도록 조심하는 방법밖에 없는 것 같다. 추가로 뒤에 연산량이 많아지는 부분 (Anti-Aliasing 구현)에서 Single/Double Precision에 따른 성능 차이를 비교해볼 계획이다.

성능 측정 결과

그림 2: CPU/GPU 실행 결과

위 그림 2는 CPU/GPU 실행 시간을 분석한 결과이다. GPU 연산 시간이 대략 2.25 배 빨라진 것을 확인 할 수 있다. 앞에서 작성한 글과 달리 연산량이 많아지면서 CPU/GPU 실행 시간 차이가 발생하는 것을 확인 할 수 있다. 뒤에 추가적인 연산을 추가하면서 성능 차이는 더 벌어질 것으로 추측된다. 혹시 코드를 원하시는 분이 있을 수도 있어서 아래 출처 4에 CPU/GPU 코드를 공유하였다.

출처

  1. https://devblogs.nvidia.com/accelerated-ray-tracing-cuda/
  2. http://www.realtimerendering.com/raytracing/Ray%20Tracing%20in%20a%20Weekend.pdf
  3. https://devblogs.nvidia.com/mixed-precision-programming-cuda-8/
  4. https://github.com/mkblog-cokr/CudaRTSeries

Leave a Comment