[CUDA RT Series 3] Adding Vectors

by

목차: Series 1 – Index + Methodology (Link)
이전 글: Series 2 –First Image (Link)
다음 글: Series 4 –Classing Up the GPU (Link)

이번 장에서는 “__global__”, “__host__”, “__device__” 사용에 관해서 설명한다. 앞에서 작성한 코드 중 GPU에서 실행될 Function에 “__global__”을 함수 이름 앞에 작성하였다. “__global__”의 경우 GPU에서 실행할 함수를 의미한다. 보통 CPU에서 처음 호출할 함수 앞에 작성한다. “__global__”로 작성된 함수를 호출하여서 연산이 수행하고 완료되는 구간을 Kernel이라고 한다. 여러 개의 “__global__” 함수를 호출하는 경우 여러 개의 Kernel 연산을 수행한 것이다. “__device__”의 경우 “__global__”과 거의 동일하게 GPU에서 연산을 수행할 함수 앞에 작성하게 된다. 하지만, “__global__” 과 달리 보통 GPU에서 호출할 함수 앞에 붙이게 된다. CPU에서 “__device__”로 작성된 함수를 호출하면 실행되지 않는다. “__host__”의 경우 CPU에서 동작할 함수 앞에 붙이게 된다. GPU에서 연산할 수 없으며 오직 CPU에서만 실행할 수 있다.

출처 2에서 Vec3 Class를 만들었다. 해당 Class는 그래픽 관련 연산을 효율적으로 하기 위해 필요한 연산을 포함한다. 이 부분을 CPU/GPU에서 모두 실행할 수 있도록 변경한다. 아래 코드 1은 Vec3 Class를 CPU/GPU에서 모두 사용할 수 있도록 변경한 코드이다.

코드 1: GPU에서 실행 가능하도록 Vec3 함수 변경 코드

#ifndef MKVEC3_H
#define MKVEC3_H
 
#include <math.h>
#include <stdlib.h>
#include <iostream>

//MK: (코드 1-1) __host__, __device__를 추가해서 CPU/GPU 모두 호출가능하도록 수정함
//MK: 그 이외 모든 부분을 기존 CPU코드와 동일함
class vec3{
    public:
      __host__ __device__ vec3(){}
      __host__ __device__ vec3(float e0, float e1, float e2){
        element[0] = e0;
        element[1] = e1;
        element[2] = e2;
      }
 
      __host__ __device__ inline float x() const{ return element[0];}
      __host__ __device__ inline float y() const{ return element[1];}
      __host__ __device__ inline float z() const{ return element[2];}
 
      __host__ __device__ inline float r() const{ return element[0];}
      __host__ __device__ inline float g() const{ return element[1];}
      __host__ __device__ inline float b() const{ return element[2];}
 
      __host__ __device__ inline const vec3& operator+() const{ return *this;}
      __host__ __device__ inline vec3 operator-() const {return vec3(-element[0], -element[1], -element[2]);}
      __host__ __device__ inline float operator[] (int i) const {return element[i];}
      __host__ __device__ inline float &operator[] (int i) {return element[i];}
 
      __host__ __device__ inline vec3& operator+=(const vec3 &v){
          element[0] += v.element[0];
          element[1] += v.element[1];
          element[2] += v.element[2];
          return *this;
      }

      __host__ __device__ inline vec3& operator-=(const vec3 &v){
          element[0] -= v.element[0];
          element[1] -= v.element[1];
          element[2] -= v.element[2];
          return *this;
      }

      __host__ __device__ inline vec3& operator*=(const vec3 &v){
          element[0] *= v.element[0];
          element[1] *= v.element[1];
          element[2] *= v.element[2];
          return *this;
      }

      __host__ __device__ inline vec3& operator/=(const vec3 &v){
          element[0] /= v.element[0];
          element[1] /= v.element[1];
          element[2] /= v.element[2];
          return *this;
      }

      __host__ __device__ inline vec3& operator*=(const float t){
          element[0] *= t;
          element[1] *= t;
          element[2] *= t;
          return *this;
      }

      __host__ __device__ inline vec3& operator/=(const float t){
          float k = 1.0/t;
          element[0] *= k;
          element[1] *= k;
          element[2] *= k;
          return *this;
      }
 
      __host__ __device__ inline float length() const{
          return sqrt(element[0] * element[0] + element[1] * element[1] + element[2] * element[2]);
      }

      __host__ __device__ inline float squared_length() const{
          return (element[0] * element[0] + element[1] * element[1] + element[2] * element[2]);
      }

	  __host__ __device__ inline void make_unit_vector(){
          float k = 1.0 / (sqrt(element[0] * element[0] + element[1] * element[1] + element[2] * element[2]));
          element[0] *= k;
          element[1] *= k;
          element[2] *= k;
      };
 
      float element[3];
};
 
inline std::istream& operator>>(std::istream &is, vec3 &t){
    is >> t.element[0] >> t.element[1] >> t.element[2];
    return is;
}
 
inline std::ostream& operator<<(std::ostream &os, const vec3 &t){
    os << t.element[0] << t.element[1] << t.element[2];
    return os;
}
 
__host__ __device__ inline vec3 operator+(const vec3 &v1, const vec3 &v2){
    return vec3(v1.element[0] + v2.element[0], v1.element[1] + v2.element[1], v1.element[2] + v2.element[2]);
}
 
__host__ __device__ inline vec3 operator-(const vec3 &v1, const vec3 &v2){
    return vec3(v1.element[0] - v2.element[0], v1.element[1] - v2.element[1], v1.element[2] - v2.element[2]);
}
 
__host__ __device__ inline vec3 operator*(const vec3 &v1, const vec3 &v2){
    return vec3(v1.element[0] * v2.element[0], v1.element[1] * v2.element[1], v1.element[2] * v2.element[2]);
}
 
__host__ __device__ inline vec3 operator/(const vec3 &v1, const vec3 &v2){
    return vec3(v1.element[0] / v2.element[0], v1.element[1] / v2.element[1], v1.element[2] / v2.element[2]);
}

__host__ __device__ inline vec3 operator*(const float t, const vec3 &v){
    return vec3(t * v.element[0], t * v.element[1], t * v.element[2]);
}
 
__host__ __device__ inline vec3 operator/(const vec3 &v, const float t){
    return vec3(v.element[0]/t, v.element[1]/t, v.element[2]/t);
}
 
__host__ __device__ inline vec3 operator*(const vec3 &v, const float t){
    return vec3(v.element[0] * t, v.element[1] * t, v.element[2] * t);
}
 
__host__ __device__ inline float dot(const vec3 &v1, const vec3 &v2){
    return (v1.element[0] * v2.element[0] + v1.element[1] * v2.element[1] + v1.element[2] * v2.element[2]);
}
 
__host__ __device__ inline vec3 cross(const vec3 &v1, const vec3 &v2){
    return vec3(
                (v1.element[1] * v2.element[2] - v1.element[2] * v2.element[1]),
                -(v1.element[0] * v2.element[2] - v1.element[2] * v2.element[0]),
                (v1.element[0] * v2.element[1] - v1.element[1] * v2.element[0])
            );
}
 
__host__ __device__ inline vec3 unitVector(vec3 v){
    return (v/v.length());
}
 
#endif

위 코드가 기존 CPU 코드와 다른 부분은 “__host__”와 “__device__”라는 문구가 모든 함수 이름 앞에 있는 것이다. “__host__”와 “__device__”를 모두 사용했기 때문에 모든 함수는 CPU와 GPU에서 실행 가능하다는 의미이다. 하지만, “__global__”이 없는 관계로 GPU에서 시작하는 함수는 아니다.

코드 2: Vec3 함수를 사용하기 위한 Main 코드 

#include <iostream>
#include <fstream>
#include <time.h>
#include "mkVec3.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 함
#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) Kernel 코드. GPU에서 수행할 함수
//MK: Vec3를 사용할 수 있도록 코드를 변경함
__global__ void mkRender(vec3 *fb, int max_x, int max_y) {
	//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 ;
	fb[pixel_index] = vec3(float(i)/max_x, float(j)/max_y, 0.2f);
}

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 = num_pixels*sizeof(vec3);

    //MK: FB 메모리 할당 (cudaMallocManaged 는 Unitifed Memory를 사용 할 수 있도록 함)
	//MK: 필요에 따라 CPU/GPU에서 GPU/CPU로 데이터를 복사함
    //MK: (코드 2-2) Vec3를 사용할 수 있도록 코드를 변경함
    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);
    checkCudaErrors(cudaGetLastError());
	//MK: CUDA 연산이 완료되길 기다림
    checkCudaErrors(cudaDeviceSynchronize());
	//MK: 연산 시간과 끝 부분을 계산하여서 연산 시간을 측정함 

    //MK: CPU 코드와 동일하게 결과를 파일에 작성
    //MK: (코드 2-3) Vec3를 사용할 수 있도록 코드를 변경함
	string fileName = "Ch3_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;
}

코드 2는 새로 작성한 Vec3 Class를 사용하기 위한 Main 코드이다. 기존 Float로 생성한 FrameBuffer (FB) 메모리를 Vec3으로 변경하였다. 추가로 “mkRender”함수 안에 Vec3 함수를 사용하는 것을 확인 할 수 있다.

그림 1: 결과 이미지

NVCC 컴파일러를 사용해서 컴파일 후 실행하면 위 그림 1을 확인 할 수 있다. 이전 글에서 작성한 코드와 동일한 결과이다. 혹시 코드를 원하시는 분이 있을 수도 있어서 아래 출처 3에 CPU/GPU 코드를 공유하였다.

 

성능 측정 결과

그림 2: CPU/GPU 실행 시간 결과

그림 2는 CPU/GPU 성능 측정 결과이다. 이전 글과 동일하게 대부분의 연산이 결과를 파일에 작성하는 부분이기 때문에 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://github.com/mkblog-cokr/CudaRTSeries

Leave a Comment