[CUDA RT Series 2] First Image

by

목차: Series 1 – Index + Methodology (Link)
이전 글: Series 1 – Index + Methodology (Link)
다음 글: Series 3 – Adding Vector (Link)

출처 1에서 가장 처음 그리는 이미지는 초록색/빨간색/파란색을 적당히 혼합한 이미지이다. 아래 그림 1은 우리가 그릴 이미지이다. 우리가 그리게 될 이미지는 크기는 1200 x 600 이다. 기존 CPU 코드는 For Loop을 여러 번(Nested For Loops) 사용해서 이미지 결과값을 계산한다. CUDA의 경우 수백/수천 개의 Thread의 연산을 동시에 수행할 수 있기 때문에 For Loop 대신 여러 개의 Thread를 생성하여서 연산을 수행한다.

그림 1: 결과 이미지

아래 코드 1은 1200 x 600 이미지를 그리기 위한 CUDA 코드이다. CUDA 코드의 확장자는 보통 “.cu”로 끝난다. 이번 글의 경우 mkMain.cu에 CUDA 코드를 작성하였다.

코드 1: mkMain.cu 코드

#include <iostream>
#include <fstream>
#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: (코드 1-1) #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: (코드 1-3) Kernel 코드. GPU에서 수행할 함수
__global__ void mkRender(float *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*3 + i*3;
    fb[pixel_index + 0] = float(i) / max_x;
    fb[pixel_index + 1] = float(j) / max_y;
    fb[pixel_index + 2] = 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 = 3*num_pixels*sizeof(float);

    //MK: (코드 1-2) FB 메모리 할당 (cudaMallocManaged 는 Unitifed Memory를 사용 할 수 있도록 함)
	//MK: 필요에 따라 CPU/GPU에서 GPU/CPU로 데이터를 복사함
    float *fb;
    checkCudaErrors(cudaMallocManaged((void **)&fb, fb_size));

    clock_t start, stop;
    start = clock();
    //MK: (코드 1-4) 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: (코드 1-5) CPU 코드와 동일하게 결과를 파일에 작성
	string fileName = "Ch1_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*3*nx + i*3;
            			float r = fb[pixel_index + 0];
            			float g = fb[pixel_index + 1];
            			float b = fb[pixel_index + 2];
            			int ir = int(255.99*r);
            			int ig = int(255.99*g);
            			int ib = int(255.99*b);
            			writeFile  << ir << " " << ig << " " << ib << "\n";
        		}
    		}
		writeFile.close();
	}

	//MK: 연산 시간과 끝 부분을 계산하여서 연산 시간을 측정함 
	stop = clock();
    double timer_seconds = ((double)(stop - start)) / CLOCKS_PER_SEC;
    cout << "MK: GPU (CUDA) Took " << timer_seconds << " Seconds.\n";
    //MK: 메모리 Free
    checkCudaErrors(cudaFree(fb));
    return 0;
}

CUDA 코드를 작성하기 이전에 먼저 에러를 확인하기 위한 코드를 작성한다. 코드 1-1 “checkCudaErrors”는 CUDA API를 사용할 때 에러가 발생하면 에러가 발생한 위치, 함수를 보여주기 위한 Macro이다. 코드 1-1은 Function 호출 후 문제가 발생하면 파일 이름, 함수 이름, 함수 위치를 표시해준다. 디버깅을 수월하게 할 수 있다. 추가로 출처 3에서 CUDA 코드 실행 시 발생하는 에러 코드에 대한 설명이 정리되어 있다.

  • MK: “__FILE__”은 Function을 호출한 파일 이름을 보여준다. “__LINE__”의 경우 에러가 발생하는 Line 위치를 의미한다. “#val”은 val을 String 값으로 리턴한다는 의미이다. 해당 Macro에서는 Function 이름을 보여준다.

이제 이미지를 그리기 위한 코드를 작성할 차례이다. 먼저 FB 값을 저장할 메모리 공간을 생성한다. 코드 1-2는 FB(1200 x 600) 메모리 공간을 생성하는 부분이다. 메모리 생성을 위해서 “cudaMallocManaged”를 사용하였다. “cudaMallocManaged”은 Unified Memory를 의미한다. 출처 1 설명에 따르면 “cudaMallocManaged”은 필요시에  CPU에서 GPU로 또는 GPU에서 CPU로 메모리값을 자동으로 복사(이동)해준다고 한다. 보통(예전)에는 cudaMalloc 함수를 사용하여 GPU 메모리 공간을 생성하고, CPU는 malloc을 사용해서 메모리 공간을 생성한다. 그리고 프로그래머가  GPU에서 CPU로 메모리값을 복사하거나, CPU에서 GPU로 메모리값을 복사하는 부분을 모두 직접 작성해야 한다 (지금도 그렇게 작성 할 수 있다). 하지만, “cudaMallocManaged”의 경우 CPU와 GPU 메모리 공간을 같이(자동으로) 사용한다는 의미인 것 같다.

다음은 GPU Thread가 실제로 연산하는 부분이다. 코드 1-3은 GPU의 Thread가 연산을 수행하기 위한 코드이다. 해당 코드는 Thread ID, Block ID, Grid Size (BlockDim) 값을 사용하여 연산할 FB의 위치 값을 계산한다. GPU에서 사용하는 Thread, Thread Block, Grid 에 대한 설명은 출처 4에 따로 작성하였다. FB 위치에 따라 값을 변경할 수 있도록 코드를 작성한다.

마지막으로 Thread Block, Grid 사이즈를 결정하여서 코드 1-3을 호출하는 부분을 작성한다. 코드 1-4는 Thread Block, Grid 사이즈를 결정하는 단계이다. 해당 코드에서는 Thread Block 사이즈를 8 x 8로 설정하였다 (dim3 threads). 다음으로 Thread Block 사이즈를 기반으로 총 생성해야 하는 Thread Block 개수 (Grid 사이즈)를 결정한다 (dim3 blocks). 이렇게 정해진 Thread Block, Grid 크기를 사용해서 코드 1-3에 작성한 함수를 호출한다. CUDA 코드의 경우 <<< >>>를 사용해서 Thread Block, Grid 크기를 알려준다.

Thread Block 사이즈를 결정하기 위해서는 여러 가지 사항을 고려해야 한다. 예를 들어 Streaming Multiprocessor (SM)의 코어 수, 실행 가능한 Thread 개수, Shared Memory (ScratchPad Memory)사이즈 등을 고려해야 한다. 가장 중요한 부분은 Thread Block 사이즈(Thread 수)를 32로 나누었을 때 나머지가 0이 되도록 (32, 64, 96, 128, …)하는 것이다.  그 이유는 GPU의 가장 기본 Execution 단위가 Warp라는 Thread 그룹 단위이기 때문이다. Warp에 대한 설명은 출처 5에 따로 작성하였다. 출처 1에서는 Thread Block 사이즈를 64로 설정하였다. 아마 여러 번 테스트를 한 결과 가장 성능이 잘 나왔기 때문으로 추측된다. 하지만 GPU Architecture마다 Thread Block 사이즈로 인한 성능 차이가 발생 할 수 있기 때문에 현재 본인이 사용하고 있는 GPU에서 적합한 사이즈는 직접 찾아야 한다.

  • MK: 뒤에 연산량이 많아지는 부분(Anti-Aliasing 구현)에서 Thread Block 사이즈를 달리하여 성능의 차이를 확인해 볼 계획이다.

이제 CUDA를 사용해서 연산을 수행한 결과값을 Output 파일로 작성한다. 코드 1-5는 GPU에서 연산한 결과 값을 파일로 쓰는 코드이다. 해당 코드는 GPU 결과값을 CPU로 복사하는 부분이 없다. 코드 1-2에서 “cudaMallocManaged”를 사용한 관계로 자동으로 GPU에서 CPU로 메모리를 복사한다.

위 코드를 NVCC 컴파일러를 사용해서 컴파일 후 실행하면 그림 1 결과 이미지를 확인 할 수 있다.

 

성능 측정 결과

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

CPU 코드(병렬화 하지 않은 코드)와 CUDA 코드(병렬화) 코드 성능을 측정하였다. 위 그림 2는 성능 결과 그래프이다. 사실상 결과의 큰 차이가 없다. 처음으로 작성한 CUDA 코드의 경우 계산 결과값을 파일에 쓰는 부분까지 시간을 측정했기 때문이다.  위 코드 1의 경우 실제 연산량은 적고 파일에 쓰는 시간이 대부분인 관계로 CPU/GPU 코드 성능차이가 거의 없다. 계속 코드를 작성하다보면 연산량이 급격히 증가하게된다. 연산량이 증가하면서 CPU/GPU 코드의 성능차이가 많이 발생하게 된다. 혹시 코드를 원하시는 분이 있을 수도 있어서 아래 출처 6에 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://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html
  4. https://mkblog.co.kr/2018/10/06/gpgpu-series-4-gpgpu-application-thread-hierarchy/
  5. https://mkblog.co.kr/2018/12/27/gpgpu-series-6-thread-block-to-warps/
  6. https://github.com/mkblog-cokr/CudaRTSeries

Leave a Comment