수업 과제 & 실습/병렬프로그래밍
GPU를 이용한 Julia Set 구현
ddongyeonn
2020. 7. 4. 15:14
1. Julia Set 정의
2. CPU 구현
#include <stdlib.h>
#include <stdio.h>
#include <math.h>
#include <time.h>
#include "..\usr\include\GL\freeglut.h"
const int Dim = 1024;
unsigned char Image[Dim * Dim * 3];
float Theta = 0.0;
int MaxIter = 256;
// 콜백 함수
void Render();
void Reshape(int w, int h);
void Timer(int id);
// 사용자 정의 함수
void CreateJuliaSet();
int Julia(float a, float b, float cx, float cy, float R);
void GetColorRainbow(float t, int &r, int &g, int &b);
int main(int argc, char **argv)
{
// GLUT 초기화
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGB);
// 윈도우 크기 설정 및 생성
glutInitWindowSize(Dim, Dim);
glutCreateWindow("Julia Set(CPU)");
// 콜백 함수 등록
glutDisplayFunc(Render);
glutReshapeFunc(Reshape);
glutTimerFunc(1, Timer, 0);
// 이벤트 처리 루프 진입
glutMainLoop();
return 0;
}
void Render()
{
// 픽셀 버퍼(배경) 흰색으로 초기화
glClearColor(1, 1, 1, 1);
glClear(GL_COLOR_BUFFER_BIT);
// Julia 집합 찾아,픽셀 버퍼를 채운다.
CreateJuliaSet();
glDrawPixels(Dim, Dim, GL_RGB, GL_UNSIGNED_BYTE, Image);
glFinish();
}
void Reshape(int w, int h)
{
glViewport(0, 0, w, h);
}
void Timer(int id)
{
// Render 함수를 호출하고, 다음 타이머를 설정한다.
Theta += 0.01;
glutPostRedisplay();
glutTimerFunc(1, Timer, 0);
}
void CreateJuliaSet()
{
float cx = cos(Theta) * 0.7885, cy = sin(Theta) * 0.7885;
float R = sqrt(cx * cx + cy * cy) * 3;
clock_t st = clock();
for (int y = 0; y < Dim; ++y){
for (int x = 0; x < Dim; ++x){
// z = a + bi;
float a = (x / (float)(Dim - 1) * 2.0 - 1.0) * R * 0.5;
float b = -(y / (float)(Dim - 1) * 2.0 - 1.0) * R * 0.5;
int offset = (y * Dim + x) * 3;
int Iter = Julia(a, b, cx, cy, R);
float t = (float)Iter / MaxIter;
int R, G, B;
GetColorRainbow(t, R, G, B);
Image[offset] = R;
Image[offset + 1] = G;
Image[offset + 2] = B;
}
}
printf("Elapsed time = %u ms\n", clock() - st);
}
int Julia(float a, float b, float cx, float cy, float R)
{
// z = a + ib;
for (int k = 0; k < MaxIter; ++k)
{
// z = z * z + c;
float re = a * a - b * b;
float im = 2 * a * b;
a = re + cx;
b = im + cy;
if (a * a + b * b > R * R)
return k;
}
return MaxIter;
}
void GetColorRainbow(float t, int &r, int &g, int &b)
{
int X = (int)(6 * t);
// 0.0 <= t <= 1.0
t = 6.0 * t - X;
switch (X)
{
case 0: // 빨강 ~ 주황
r = 255;
g = t * 127;
b = 0;
//r = g = b = 255;
break;
case 1: // 주황 ~ 노랑
r = 0;
g = (1.0 - t) * 127 + t * 255;
b = 0;
break;
case 2: // 노랑 ~ 초록
r = (1.0 - t) * 255;
g = 255;
b = 0;
break;
case 3: // 초록 ~ 파랑
r = 0;
g = (1.0 - t) * 255;
b = t * 255;
break;
case 4: // 파랑 ~ 남색
r = t * 75;
g = 0;
b = (1.0 - t) * 255 + t * 130;
break;
case 5: // 남색 ~ 보라
r = (1.0 - t) * 75 + t * 148;
g = 0;
b = (1.0 - t) * 130 + t * 211;
break;
case 6: // 보라
r = 148;
g = 0;
b = 211;
break;
}
// 범위 밖은 검정색을 할당한다.
if (t < 0.0 || t > 1.0)
r = g = b = 0;
}
-
CreateJuliaSet
-
지정한 범위 이내의 복소수 값을 할당 -> julia연산 진행 -> 픽셀별 색상을 결정하는 함수
-
임의의 복소수 C의 실수부와 허수부, 그리고 복소수를 할당할 범위 R은 변하는 Theta값을 전달받은 Sin, Cos 함수에 의해 지속적으로 바뀌게 설정한다.
-
설정된 DIM*DIM 개 만큼 |R| 범위 이내에서 복소수 z를 생성한다. 즉 전체 픽셀이 DIM*DIM개 이므로 각 픽셀별로 범위내의 고유 복소수를 할당하게 되고 Julia 점화식을 진행하여 픽셀의 색상을 결정한다.
-
할당한 복소수 z값을 Julia 함수에 전달하여 점화식을 계산 후 계산 결과를 GetColorRainbow 함수에 픽셀의 전달하여 색상을 지정한다.
-
-
Julia
- 전달받은 복소수 z를 이용하여 julia연산을 진행하고 반환값으로 점화식이 진행된 횟수인 iter값을 전달하는 역할을 함
-
GetColorRainbow
- 픽셀별로 할당된 복소수 z를 이용해 julia연산을 진행하면 iter값이 반환되는데, 이 값을 통해 픽셀의 색상을 결정하는 함수
3. GPU 구현
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <time.h>
#include <memory.h>
#include "..\usr\include\GL\freeglut.h"
const int Dim = 1024;
const int MaxIter = 256; //최대 반복횟수를 256
const int TILE_WIDTH = 32;
float Theta = 0.0;
unsigned char pHostImage[Dim * Dim * 3]; //1024 * 1024 이미지 배열 -> RGB성분(x3)
// 콜백 함수
void Render();
void Reshape(int w, int h);
void Timer(int id);
// 사용자 정의 함수
void CreateThread(); // 이미지배열 만들어내는 함수
__device__ int Julia(float a, float b, float cx, float cy, float R); // 어떤 복소수가 JULIASET에 포함되는지 판단하는 함수
__device__ void GetColorRainbow(float t, int& r, int& g, int& b); // 무지개색을 얻어내는 함수
__global__ void CreateJuliaSetGPU(unsigned char* pDeviceImage, float cx, float cy, float R);
int main(int argc, char** argv)
{
// GLUT 초기화
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGB);
// 윈도우 크기 설정 및 생성
glutInitWindowSize(Dim, Dim);
glutCreateWindow("Julia Set(GPU)");
// 콜백 함수 등록
glutDisplayFunc(Render);
glutReshapeFunc(Reshape);
glutTimerFunc(30, Timer, 0); //1초에 30번 화면 호출
// 이벤트 처리 루프 진입
glutMainLoop();
return 0;
}
void Render()
{
// 픽셀 버퍼(배경) 흰색으로 초기화
glClearColor(1, 1, 1, 1);
glClear(GL_COLOR_BUFFER_BIT);
CreateThread();
glDrawPixels(Dim, Dim, GL_RGB, GL_UNSIGNED_BYTE, pHostImage);
glFinish();
}
void Reshape(int w, int h)
{
glViewport(0, 0, w, h);
}
void Timer(int id)
{
// Render 함수를 호출하고, 다음 타이머를 설정한다.
Theta += 0.01;
glutPostRedisplay();
glutTimerFunc(30, Timer, 0); //30ms마다 Timer함수 호출
}
void CreateThread() {
// C = cx + cyi 값 설정 -> 세타 값에 따라 계속 변경
float cx = cos(Theta) * 0.7885;
float cy = sin(Theta) * 0.7885;
float R = sqrt(cx * cx + cy * cy) * 3; // 줄리아 셋 연산을 위한 특정 상수 R 설정
// GPU 이미지 배열 생성 및 할당
unsigned char* pDeviceImage = NULL;
cudaError_t cudaStatus = cudaSetDevice(0);
cudaMalloc((void**)&pDeviceImage, Dim * Dim * 3 * sizeof(unsigned char));
// grid 및 block 생성
dim3 dimGrid(Dim / TILE_WIDTH, Dim / TILE_WIDTH); // 그리드 안 32 * 32 블록이 생성
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH); // 블록 안 32 * 32 스레드 생성
// 커널함수 호출 - Julia 집합
clock_t st = clock();
CreateJuliaSetGPU << <dimGrid, dimBlock >> > (pDeviceImage, cx, cy, R);
printf("Elapsed time = %u ms\n", clock() - st);
cudaDeviceSynchronize();
//결과 전송
cudaMemcpy(pHostImage, pDeviceImage, sizeof(unsigned char) * Dim * Dim * 3, cudaMemcpyDeviceToHost);
// 메모리 해제
cudaFree(pDeviceImage);
}
__global__ void CreateJuliaSetGPU(unsigned char* pDeviceImage, float cx, float cy, float R) {
//스레드의 번호 부여
int j = blockIdx.y * TILE_WIDTH + threadIdx.y; // 스레드의 행 번호
int i = blockIdx.x * TILE_WIDTH + threadIdx.x; // 스레드의 열 번호
// 현재 픽셀 (i,j) 를 z0 = a + bi로 각각 x, y 좌표 -0.5R~ +0.5R 범위의 a + bi로 변환
float a = (i / (float)(Dim - 1) * 2.0 - 1.0) * R * 0.5;
float b = -(j / (float)(Dim - 1) * 2.0 - 1.0) * R * 0.5;
// Julia Set 포함여부 확인
int Iter = Julia(a, b, cx, cy, R); // 반환된 가능한 n 까지 Iter에 저장
float t = (float)Iter / MaxIter; // 반복 횟수에 따라 t를 0~1 까지 변환
// t값을(0~1) 빨주노초파남보로 매핑 (많이반복할수록 색이 달라짐)
int Red, Green, Blue;
GetColorRainbow(t, Red, Green, Blue);
// 1차원 배열 idx값으로 변환
int offset = (j * Dim + i) * 3; //시작 idx
pDeviceImage[offset] = Red;
pDeviceImage[offset + 1] = Green;
pDeviceImage[offset + 2] = Blue;
}
__device__ int Julia(float a, float b, float cx, float cy, float R)
{
// z0 = a + bi; C = cx + cy*i; R은 최대 범위
// MaxIter -> zn 과정을 몇번 할지 (=256)
for (int k = 0; k < MaxIter; ++k)
{
// z(n+1) = z(n) * z(n) + c;
float re = a * a - b * b; //실수부
float im = 2 * a * b; //허수부
a = re + cx;
b = im + cy;
// 만약 R제곱 보다 크다면 JULIA SET 에 포함 X -> 현재까지 반복 수 return
if (a * a + b * b > R * R)
return k;
}
//최대 반복해도 julia set에 포함
return MaxIter;
}
__device__ void GetColorRainbow(float t, int& r, int& g, int& b)
{
int X = (int)(6 * t);
// 0.0 <= t <= 1.0
t = 6.0 * t - X;
switch (X)
{
case 0: // 빨강 ~ 주황
r = 255;
g = t * 127;
b = 0;
//r = g = b = 255;
break;
case 1: // 주황 ~ 노랑
r = 0;
g = (1.0 - t) * 127 + t * 255;
b = 0;
break;
case 2: // 노랑 ~ 초록
r = (1.0 - t) * 255;
g = 255;
b = 0;
break;
case 3: // 초록 ~ 파랑
r = 0;
g = (1.0 - t) * 255;
b = t * 255;
break;
case 4: // 파랑 ~ 남색
r = t * 75;
g = 0;
b = (1.0 - t) * 255 + t * 130;
break;
case 5: // 남색 ~ 보라
r = (1.0 - t) * 75 + t * 148;
g = 0;
b = (1.0 - t) * 130 + t * 211;
break;
case 6: // 보라
r = 148;
g = 0;
b = 211;
break;
}
// 범위 밖은 검정색을 할당한다.
if (t < 0.0 || t > 1.0)
r = g = b = 0;
}
-
CPU 버전과 기본적인 구조는 동일하지만 복소수를 할당하는 부분에서 차이가 있다.
-
CPU 버전에는 2중 for문을 통해 복소수를 할당한 반면 GPU를 이용하게 되면 커널함수를 통해 DIM*DIM 개의 스레드를 생성하게 되고 각 스레드가 복소수를 할당 -> Julia 연산 -> 색상지정의 과정을 병렬적으로 실행하게 된다.
-
CreateThread() 함수에서 임의복소수 C와 범위 R을 설정한 뒤에 스레드 관리를 위한 Grid와 Block을 설정한다.
-
CreateJuliaSetGPU() 커널함수를 호출하게 되면 전달된 Grid, Block정보를 이용해 스레드를 생성하고 각 스레드가 앞서 설명한 과정을 병렬적으로 진행한다.
-
CudaDeviceSynchronize() 함수를 이용해 모든 스레드를 동기화 해준 뒤 결과값을 CPU로 전달해 화면을 출력한다.
4. 결과영상