ddongstudy
Newton Fractal 본문
1. Newton Fractal 정의
2. CPU 구현
// CPU(점&픽셀) 최종
#include "..\usr\include\GL\freeglut.h"
#include <math.h>
#include <stdio.h>
#include <time.h>
#include<stdlib.h>
#include<complex>
#include<memory.h>
struct ucomplex {
float real;
float imag;
};
const int width = 1024;
float scope = 0.8; // 확대,축소
float pre_scope = 0.0;
float* z = (float*)malloc(width * width * sizeof(float));
unsigned char Image[width * width * 3];
float dx = 0.0, dy = 0.0; // mouse move
bool TimeDir = true, Mode = true; // 확대 축소 방향 . time/scroll
int key_s = 0; // 방정식 바꾸는 key
int key_d = 0; // 차수 늘리는 key
float Zoom = -50.0;
int StartPt[2];
// 콜백 함수
void Keyboard(unsigned char key, int x, int y);
void Mouse(int button, int state, int x, int y);
void MouseMove(int x, int y);
void MouseWheel(int button, int dir, int x, int y);
void Reshape(int w, int h);
void Timer(int id);
// 연산자 오버라이딩
////////////////////////////////////////////////////////////////////////////////////////
float abs(ucomplex x) {
return sqrt(x.real * x.real + x.imag * x.imag);
}
ucomplex operator*(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real * y.real - x.imag * y.imag),
(x.imag * y.real + x.real * y.imag)
};
return c;
}
ucomplex operator*(ucomplex x, float dy) {
ucomplex y = { dy, 0.0 };
return x * y;
}
ucomplex operator/(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real * y.real + x.imag * y.imag) / (y.real * y.real + y.imag * y.imag),
(x.imag * y.real - x.real * y.imag) / (y.real * y.real + y.imag * y.imag)
};
return c;
}
ucomplex operator+(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real + y.real),
(x.imag + y.imag)
};
return c;
}
ucomplex operator+(ucomplex x, float dy) {
ucomplex y = { dy, 0.0 };
return x + y;
}
ucomplex operator-(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real - y.real),
(x.imag - y.imag)
};
return c;
}
ucomplex operator-(ucomplex x, float dy) {
ucomplex y = { dy, 0.0 };
return x - y;
}
/////////////////////////////////////////////////////////////////////////////////////////
ucomplex f(ucomplex x)
{
// 1. 초기 방정식: x^3 - 1
ucomplex d = x * x * x - 1.0;
if (key_s == 0)
{
ucomplex temp_d = x * x * x;
if (key_d > 0) {
for (int i = 0; i < key_d; i++)
temp_d = temp_d * x;
}
d = temp_d - 1.0;
}
// 2. x^3 - 2x^2 + 2
else if (key_s == 1)
{
d = x * x * x - x * x * 2.0 + 2.0;
}
// 3. x^8 + 15x^4 - 16
else if (key_s == 2)
{
d = x * x * x * x * x * x * x * x + x * x * x * x * 15.0 - 16.0;
}
// 4. sin(x)
else if (key_s == 3)
{
typedef std::complex<float> dcomplex;
dcomplex a(x.real, x.imag);
d = {
sin(a).real(),
sin(a).imag()
};
}
return d;
}
ucomplex df(ucomplex x) // f 미분값
{
ucomplex d = x * x * 3;
// 1. 3x^2
if (key_s == 0)
{
ucomplex temp_d = x * x;
float de = 3.0;
if (key_d > 0) {
for (int i = 0; i < key_d; i++) {
temp_d = temp_d * x;
de = de + 1.0;
}
}
d = temp_d * de;
}
// 2. 3x^2 - 4x
else if (key_s == 1)
{
d = x * x * 3 - x * 4.0;
}
// 3. 8x^7 +60x^3
else if (key_s == 2)
{
d = x * x * x * x * x * x * x * 8.0 + x * x * x * 60.0;
}
// 4. cos(x)
else if (key_s == 3)
{
typedef std::complex<float> dcomplex;
dcomplex a(x.real, x.imag);
d = {
cos(a).real(),
cos(a).imag()
};
}
return d;
}
//사용자 정의함수
__int64 GetMicroSecond();
float newton(ucomplex x0, float eps, int maxiter);
void MathCPU();
void Display();
int main(int argc, char** argv) {
// opengl 초기화
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGB);
//glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGB);
// 윈도우 설정
glutInitWindowSize(width, width);
//glutInitWindowSize(700, 700);
glutInitWindowPosition(0, 0);
glutCreateWindow("Newton");
glClearColor(1.0, 1.0, 1.0, 1.0);
// 관측설정
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
glOrtho(-2, 2, 2, -2, 1, -1);
// 콜백함수 등록
glutDisplayFunc(Display);
glutTimerFunc(30, Timer, 0);
glutReshapeFunc(Reshape);
glutKeyboardFunc(Keyboard);
glutMouseFunc(Mouse);
glutMotionFunc(MouseMove);
glutMouseWheelFunc(MouseWheel);
glutMainLoop();
}
void Keyboard(unsigned char key, int x, int y)
{
if (key == 27)
exit(-1);
// 방정식 바꾸기
if (key == 's')
{
key_d = 0;
key_s = (key_s + 1) % 4;
}
// 차수 늘리기
if (key == 'd')
{
key_d++;
}
if (key == 'e')
{
if (key_d > 0)
key_d--;
}
// 모드 바꾸기
if (key == 'f') {
Mode = !Mode;
}
glutPostRedisplay();
}
void Mouse(int button, int state, int x, int y)
{
// 마우스 버튼을 누른 경우,
if (state == GLUT_DOWN)
{
StartPt[0] = x;
StartPt[1] = y;
}
// 마우스 버튼을 땐 경우,
if (state == GLUT_UP)
StartPt[0] = StartPt[1] = 0;
}
void MouseMove(int x, int y) {
dx += -(x - StartPt[0]) * 0.01; // x 이동값
dy += -(StartPt[1] - y) * 0.01; // y 이동값
}
void MouseWheel(int button, int dir, int x, int y)
{
// Mode = false 일때
if (!Mode) {
if (dir > 0)
{
pre_scope = scope;
scope -= 0.01;
if (scope <= 0.00001)
scope = 0.00001;
}
else
{
pre_scope = scope;
scope += 0.01;
if (scope >= 1.57)
scope = 1.57;
}
}
glutPostRedisplay();
}
float newton(ucomplex x0, float eps, int maxiter) {
ucomplex x = x0;
int iter = 0;
// 픽셀의 고유 iter값 계산
while (abs(f(x)) > eps && iter <= maxiter) {
iter++;
x = x - f(x) / df(x); // 점화식
}
return iter;
}
void Timer(int id)
{
// Mode = true 일때 시간에따라 확대,축소
if (Mode) {
if (TimeDir == true) {
scope -= 0.05;
if (scope <= 0.00001) {
scope = 0.00001;
TimeDir = false;
}
}
else {
scope += 0.05;
if (scope >= 1.57) {
scope = 1.57;
TimeDir = true;
}
}
}
glutPostRedisplay();
glutTimerFunc(30, Timer, 0); //30ms마다 Timer함수 호출
}
void Reshape(int w, int h)
{
glViewport(0, 0, w, h);
}
void MathCPU() {
// 범위 설정
float xmin = -2, xmax = 2;
float ymin = -2, ymax = 2;
// 각 픽셀간의 간격을 설정
int xsteps = width, ysteps = width;
float hx = (xmax - xmin) / xsteps, hy = (ymax - ymin) / ysteps;
float eps = 0.0001; // eps보다 작아지면 더이상 진행하지 않음
int maxiter = 255; // 최대 iter값
float scopenum = 15 * sin(scope); // 확대 축소 범위: -15 ~ 15
float x, y;
y = ymin;
// 픽셀별 뉴턴값 계산
for (int i = 0; i < ysteps; i++) {
x = xmin;
for (int j = 0; j < xsteps; j++) {
ucomplex xy = { x * scopenum + dx,y * scopenum + dy }; // 복소수값 할당
z[i * width + j] = newton(xy, eps, maxiter); // 뉴턴값 계산
x += hx; // + 1열
}
y += hy; // +1행
}
}
// width * width 픽셀에 색상을 지정
void Display() {
glClearColor(1, 1, 1, 1);
glClear(GL_COLOR_BUFFER_BIT);
__int64 _st = GetMicroSecond();
MathCPU();
for (int i = 0; i < width; i++) {
for (int j = 0; j < width; j++) {
float color = (z[i * width + j] / (float)256);
int offset = (i * width + j) * 3;
Image[offset] = fmod(color * 13, 1.0f) * 255;
Image[offset + 1] = fmod(color * 33, 1.0f) * 255;
Image[offset + 2] = fmod(color * 49, 1.0f) * 255;
}
}
glDrawPixels(width, width, GL_RGB, GL_UNSIGNED_BYTE, Image);
printf("Elapsed time: %I64d mico sec \n", GetMicroSecond() - _st);
glFinish();
}
__int64 GetMicroSecond()
{
LARGE_INTEGER frequency;
LARGE_INTEGER now;
if (!QueryPerformanceFrequency(&frequency))
return (__int64)GetTickCount();
if (!QueryPerformanceCounter(&now))
return (__int64)GetTickCount();
return((now.QuadPart) / (frequency.QuadPart / 1000000));
}
/*
//width * width 점을 찍어서 렌더링
void Display() {
glClear(GL_COLOR_BUFFER_BIT);
glBegin(GL_POINTS);
__int64 _st = GetMicroSecond();
MathCPU();
float xmin = -2, xmax = 2;
float ymin = -2, ymax = 2;
int xsteps = width, ysteps = width;
float hx = (xmax - xmin) / xsteps, hy = (ymax - ymin) / ysteps;
float x, y;
float max = z[0];
float min = z[0];
for (int i = 0; i < width; i++) {
for (int j = 0; j < width; j++) {
if (z[i * width + j] >= max) max = z[i * width + j];
if (z[i * width + j] <= min) min = z[i * width + j];
}
}
y = ymin;
for (int i = 0; i < width; i++) {
x = xmin;
for (int j = 0; j < width; j++) {
float color = (z[i * width + j]) / (256);
glColor3d(fmod(color * 13, 1.0f),
fmod(color * 33, 1.0f),
fmod(color * 49, 1.0f));
glVertex2d(x, y);
x += hx;
}
y += hy;
}
printf("Elapsed time: %I64d mico sec \n", GetMicroSecond() - _st);
glEnd();
glutSwapBuffers();
}
*/-
newton()
-
임의의 복소수값을 전달받으면 newton점화식을 수렴할때까지 반복하고 반복된 횟수를 반환한다.
-
-
MathCPU()
-
픽셀별로 복소수를 할당한다.
-
복소수를 할당할 x, y범위를 설정하고 범위 내에서 픽셀의 개수(width*width)만큼 복소수 값을 설정한다.
-
생성된 복소수 별로 newton점화식을 계산하고 결과값을 z배열에 저장한다.
-
-
Display()
-
z배열에 저장된 픽셀별 newton점화식 값을 이용해 픽셀의 색상을 결정한다.
-
3. GPU 구현
// GPU - PBO 이용
#include <math.h>
//#include <glut.h>
#include <stdio.h>
#include <time.h>
#include <cuda.h>
#include <iostream>
#include <complex>
#include <cuComplex.h>
#include <thrust/complex.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "..\usr\include\GL\glew.h"
#include "..\usr\include\GL\freeglut.h"
#include "cuda_gl_interop.h"
using namespace std;
// 복소수 구조체
struct ucomplex {
float real;
float imag;
};
const int width = 1024; // 넓이
float scope = 0.8; // 확대,축소
float pre_scope = 0.0;
float Dx = 0.0, Dy = 0.0; // mouse move
bool TimeDir = true, Mode = true; // 확대 축소 방향 . time/scroll
int key_s = 0; // 방정식 바꾸는 key
int key_d = 0; // 차수 늘리는 key
float Zoom = -50.0;
int StartPt[2];
//GPU 관련 변수
dim3 dimGrid;
dim3 dimBlock;
float* zD; // 뉴턴방정식 계산 결과를 저장할 배열
float p = 0;
GLuint gl_pbo; // 픽셀 버퍼를 가르키는 OpenGL 핸들
cudaGraphicsResource* cuda_pbo; // 픽셀 버퍼를 가르키는 CUDA 핸들
uchar4* pDevImage; // 픽셀 버퍼에 대한 실제 메모리 주소
// 콜백 함수
void Keyboard(unsigned char key, int x, int y);
void Mouse(int button, int state, int x, int y);
void MouseMove(int x, int y);
void MouseWheel(int button, int dir, int x, int y);
void Reshape(int w, int h);
void Timer(int id);
void Close();
void Display();
//사용자 정의함수
__global__ void MathGPUKernel(float* zD, float scopenum, float dx, float dy, int Key_d, int Key_s, uchar4* DevImage);
__int64 GetMicroSecond();
__host__ __device__ float abs(ucomplex x) {
return sqrt(x.real * x.real + x.imag * x.imag);
}
__host__ __device__ ucomplex operator*(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real * y.real - x.imag * y.imag),
(x.imag * y.real + x.real * y.imag)
};
return c;
}
__host__ __device__ ucomplex operator*(ucomplex x, float dy) {
ucomplex y = { dy, 0.0 };
return x * y;
}
__host__ __device__ ucomplex operator/(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real * y.real + x.imag * y.imag) / (y.real * y.real + y.imag * y.imag),
(x.imag * y.real - x.real * y.imag) / (y.real * y.real + y.imag * y.imag)
};
return c;
}
__host__ __device__ ucomplex operator+(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real + y.real),
(x.imag + y.imag)
};
return c;
}
__host__ __device__ ucomplex operator+(ucomplex x, float dy) {
ucomplex y = { dy, 0.0 };
return x + y;
}
__host__ __device__ ucomplex operator-(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real - y.real),
(x.imag - y.imag)
};
return c;
}
__host__ __device__ ucomplex operator-(ucomplex x, float dy) {
ucomplex y = { dy, 0.0 };
return x - y;
}
__host__ __device__ ucomplex f(ucomplex x, int Key_d, int Key_s) {
// 1. 초기 방정식: x^3 - 1
ucomplex d = x * x * x - 1.0;
if (Key_s == 0)
{
ucomplex temp_d = x * x * x;
if (Key_d > 0) {
for (int i = 0; i < Key_d; i++)
temp_d = temp_d * x;
}
d = temp_d - 1.0;
}
// 2. x^3 - 2x^2 + 2
else if (Key_s == 1)
{
d = x * x * x - x * x * 2.0 + 2.0;
}
// 3. x^8 + 15x^4 - 16
else if (Key_s == 2)
{
d = x * x * x * x * x * x * x * x + x * x * x * x * 15.0 - 16.0;
}
// 4. sin(x)
else if (Key_s == 3)
{
thrust::complex<float> c = thrust::complex<float>(x.real, x.imag);
d = {
sin(c).real(),
sin(c).imag()
};
}
return d;
}
__host__ __device__ ucomplex df(ucomplex x, int Key_d, int Key_s) {
ucomplex d = x * x * 3;
// 1. 3x^2
if (Key_s == 0)
{
ucomplex temp_d = x * x;
float de = 3.0;
if (Key_d > 0) {
for (int i = 0; i < Key_d; i++) {
temp_d = temp_d * x;
de = de + 1.0;
}
}
d = temp_d * de;
}
// 2. 3x^2 - 4x
else if (Key_s == 1)
{
d = x * x * 3 - x * 4.0;
}
// 3. 8x^7 +60x^3
else if (Key_s == 2)
{
d = x * x * x * x * x * x * x * 8.0 + x * x * x * 60.0;
}
// 4. cos(x)
else if (Key_s == 3)
{
thrust::complex<float> c = thrust::complex<float>(x.real, x.imag);
d = {
cos(c).real(),
cos(c).imag()
};
}
return d;
}
__host__ __device__ float newton(ucomplex x0, float eps, int maxiter, int Key_d, int Key_s) {
ucomplex x = x0;
int iter = 0;
// 뉴턴식 계산 - 픽셀의 고유 iter값을 얻는다
while (abs(f(x, Key_d, Key_s)) > eps && iter <= maxiter) {
iter++;
x = x - f(x, Key_d, Key_s) / df(x, Key_d, Key_s); // 점화식
}
return iter;
}
int main(int argc, char** argv) {
// opengl 초기화
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGBA);
//glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGB);
// 윈도우 설정
glutInitWindowSize(width, width);
//glutInitWindowSize(700, 700);
glutInitWindowPosition(0, 0);
glutCreateWindow("Newton");
// 관측설정
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
glOrtho(-2, 2, 2, -2, 1, -1);
// 콜백함수 등록
glutDisplayFunc(Display);
glutTimerFunc(30, Timer, 0);
glutReshapeFunc(Reshape);
glutKeyboardFunc(Keyboard);
glutMouseFunc(Mouse);
glutMotionFunc(MouseMove);
glutMouseWheelFunc(MouseWheel);
glutCloseFunc(Close);
// grid, block 만들기
const int block_width = 16;
cudaMalloc((void**)&zD, width * width * sizeof(float));
dimGrid = dim3((width - 1 / block_width) + 1, (width - 1 / block_width) + 1);
dimBlock = dim3(block_width, block_width);
// glew 초기화
glewInit();
// GPU Device 세팅
cudaSetDevice(0);
cudaGLSetGLDevice(0);
glGenBuffers(1, &gl_pbo);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, gl_pbo);
// 픽셀 버퍼 할당
glBufferData(GL_PIXEL_UNPACK_BUFFER_ARB, width * width * sizeof(uchar4), NULL, GL_DYNAMIC_DRAW_ARB);
// 이벤트 처리 루프 진입
glutMainLoop();
// 동적할당 해제
cudaFree(zD);
}
void MathGPU() {
// 화면 비율에 맞춰 확대축소 비율 계산
if (width < 1000)
p = 1;
else
p = (width / 999);
float scopenum = p * 15 * sin(scope);
float dx = Dx, dy = Dy;
MathGPUKernel << <dimGrid, dimBlock >> > (zD, scopenum, dx, dy, key_d, key_s, pDevImage); // 커널함수 호출
cudaDeviceSynchronize(); // 동기화
}
__global__ void MathGPUKernel(float* zD, float scopenum, float dx, float dy, int Key_d, int Key_s, uchar4* DevImage) {
float xi = blockIdx.x * blockDim.x + threadIdx.x;
float yi = blockIdx.y * blockDim.y + threadIdx.y;
int maxiter = 255;
float eps = 0.0001;
// 뉴턴 점화식 계산
uchar4 C;
if (xi < width && yi < width) {
// 각 픽셀별 복소수값 할당 후 뉴턴값 계산
// 마우스 스크롤(scope) -> 전체 범위가 확대 축소(곱 연산)
// 마우스 드래그(dx,dy) -> 범위의 크기는 그대로이고 범위가 평행이동(+연산)
ucomplex xy = { (xi / width * 4 - 2) * scopenum + dx, (yi / width * 4 - 2) * scopenum + dy };
zD[(int)yi * width + (int)xi] = newton(xy, eps, maxiter, Key_d, Key_s);
// 색상 계산
float color = zD[(int)yi * width + (int)xi] / maxiter;
int offset = ((int)yi * width + (int)xi);
C.x = color * 13 * 255;
C.y = color * 33 * 255;
C.z = color * 49 * 255;
C.w = 255;
DevImage[offset] = C;
}
}
// width * width 픽셀에 색상을 지정
void Display() {
// 배경 초기화
__int64 _st = GetMicroSecond();
glClearColor(1, 1, 1, 1);
glClear(GL_COLOR_BUFFER_BIT);
// 픽셀 버퍼를 CUDA에 등록
cudaGraphicsGLRegisterBuffer(&cuda_pbo, gl_pbo, cudaGraphicsMapFlagsNone);
// 픽셀 버퍼를 CUDA 시스템에 mapping
size_t size;
cudaGraphicsMapResources(1, &cuda_pbo, NULL);
cudaGraphicsResourceGetMappedPointer((void**)&pDevImage, &size, cuda_pbo);
MathGPU();
glDrawPixels(width, width, GL_RGBA, GL_UNSIGNED_BYTE, 0);
printf("Elapsed time: %I64d mico sec \n", GetMicroSecond() - _st);
glFinish();
cudaGraphicsUnmapResources(1, &cuda_pbo, NULL);
cudaGraphicsUnregisterResource(cuda_pbo);
}
void Close() {
// GPU 리셋
cudaDeviceReset();
//픽셀 버퍼의 바인딩 해제 및 핸들 제거
glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
glDeleteBuffers(1, &gl_pbo);
}
void Timer(int id)
{
// Mode = true 일때 시간에따라 자동 확대, 축소 (key = f)
// TimeDir = true 이면 축소 false이면 확대
if (Mode) {
if (TimeDir == true) {
scope -= 0.05;
if (scope <= 0.00001) {
scope = 0.00001;
TimeDir = false;
}
}
else {
scope += 0.05;
if (scope >= 1.57) {
scope = 1.57;
TimeDir = true;
}
}
}
glutPostRedisplay();
glutTimerFunc(30, Timer, 0); //30ms마다 Timer함수 호출
}
void Reshape(int w, int h)
{
glViewport(0, 0, w, h);
}
void Keyboard(unsigned char key, int x, int y)
{
if (key == 27)
exit(-1);
// 방정식 바꾸기
if (key == 's')
{
key_d = 0;
key_s = (key_s + 1) % 4;
}
// 차수 늘리기
if (key == 'd')
{
key_d++;
}
// 차수 낮추기
if (key == 'e')
{
if (key_d > 0)
key_d--;
}
// 자동 확대 모드 바꾸기
if (key == 'f') {
Mode = !Mode;
}
glutPostRedisplay();
}
void Mouse(int button, int state, int x, int y)
{
// 마우스 버튼을 누른 경우,
if (state == GLUT_DOWN)
{
StartPt[0] = x;
StartPt[1] = y;
}
// 마우스 버튼을 땐 경우,
if (state == GLUT_UP)
StartPt[0] = StartPt[1] = 0;
}
void MouseMove(int x, int y) {
Dx += -(x - StartPt[0]) * 0.01; // x 이동값
Dy += -(StartPt[1] - y) * 0.01; // y 이동값
}
void MouseWheel(int button, int dir, int x, int y)
{
// Mode = false 일때
// 수동 확대 축소
if (!Mode) {
if (dir > 0)
{
pre_scope = scope;
scope -= 0.01;
if (scope <= 0.00001)
scope = 0.00001;
}
else
{
pre_scope = scope;
scope += 0.01;
if (scope >= 1.57)
scope = 1.57;
}
}
glutPostRedisplay();
}
__int64 GetMicroSecond()
{
LARGE_INTEGER frequency;
LARGE_INTEGER now;
if (!QueryPerformanceFrequency(&frequency))
return (__int64)GetTickCount();
if (!QueryPerformanceCounter(&now))
return (__int64)GetTickCount();
return((now.QuadPart) / (frequency.QuadPart / 1000000));
}
// GPU - 픽셀에 색상을 지정
/*
#include "..\usr\include\GL\freeglut.h"
#include <math.h>
//#include <glut.h>
#include <stdio.h>
#include <time.h>
#include <cuda.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include <complex>
#include <cuComplex.h>
#include <thrust/complex.h>
using namespace std;
struct ucomplex {
float real;
float imag;
};
const int width = 1024;
float scope = 0.8; // 확대,축소
float pre_scope = 0.0;
float* z = (float*)malloc(width * width * sizeof(float));
unsigned char Image[width * width * 3];
float Dx = 0.0, Dy = 0.0; // mouse move
bool TimeDir = true, Mode = true; // 확대 축소 방향 . time/scroll
int key_s = 0; // 방정식 바꾸는 key
int key_d = 0; // 차수 늘리는 key
float Zoom = -50.0;
int StartPt[2];
//GPU 관련 변수
dim3 dimGrid;
dim3 dimBlock;
float* zD;
unsigned char* DevImage;
int Size;
float p = 0;
// 콜백 함수
void Keyboard(unsigned char key, int x, int y);
void Mouse(int button, int state, int x, int y);
void MouseMove(int x, int y);
void MouseWheel(int button, int dir, int x, int y);
void Reshape(int w, int h);
void Timer(int id);
//사용자 정의함수
void MathCPU();
void Display();
__int64 GetMicroSecond();
__global__ void MathGPUKernel(float* zD, float scopenum, float dx, float dy, int Key_d, int Key_s, unsigned char* DevImage);
__host__ __device__ float abs(ucomplex x) {
return sqrt(x.real * x.real + x.imag * x.imag);
}
__host__ __device__ ucomplex operator*(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real * y.real - x.imag * y.imag),
(x.imag * y.real + x.real * y.imag)
};
return c;
}
__host__ __device__ ucomplex operator*(ucomplex x, float dy) {
ucomplex y = { dy, 0.0 };
return x * y;
}
__host__ __device__ ucomplex operator/(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real * y.real + x.imag * y.imag) / (y.real * y.real + y.imag * y.imag),
(x.imag * y.real - x.real * y.imag) / (y.real * y.real + y.imag * y.imag)
};
return c;
}
__host__ __device__ ucomplex operator+(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real + y.real),
(x.imag + y.imag)
};
return c;
}
__host__ __device__ ucomplex operator+(ucomplex x, float dy) {
ucomplex y = { dy, 0.0 };
return x + y;
}
__host__ __device__ ucomplex operator-(ucomplex x, ucomplex y) {
ucomplex c = {
(x.real - y.real),
(x.imag - y.imag)
};
return c;
}
__host__ __device__ ucomplex operator-(ucomplex x, float dy) {
ucomplex y = { dy, 0.0 };
return x - y;
}
__host__ __device__ ucomplex f(ucomplex x, int Key_d, int Key_s) {
// 1. 초기 방정식: x^3 - 1
ucomplex d = x * x * x - 1.0;
if (Key_s == 0)
{
ucomplex temp_d = x * x * x;
if (Key_d > 0) {
for (int i = 0; i < Key_d; i++)
temp_d = temp_d * x;
}
d = temp_d - 1.0;
}
// 2. x^3 - 2x^2 + 2
else if (Key_s == 1)
{
d = x * x * x - x * x * 2.0 + 2.0;
}
// 3. x^8 + 15x^4 - 16
else if (Key_s == 2)
{
d = x * x * x * x * x * x * x * x + x * x * x * x * 15.0 - 16.0;
}
// 4. sin(x)
else if (Key_s == 3)
{
thrust::complex<float> c = thrust::complex<float>(x.real, x.imag);
d = {
sin(c).real(),
sin(c).imag()
};
}
return d;
}
__host__ __device__ ucomplex df(ucomplex x, int Key_d, int Key_s) {
ucomplex d = x * x * 3;
// 1. 3x^2
if (Key_s == 0)
{
ucomplex temp_d = x * x;
float de = 3.0;
if (Key_d > 0) {
for (int i = 0; i < Key_d; i++) {
temp_d = temp_d * x;
de = de + 1.0;
}
}
d = temp_d * de;
}
// 2. 3x^2 - 4x
else if (Key_s == 1)
{
d = x * x * 3 - x * 4.0;
}
// 3. 8x^7 +60x^3
else if (Key_s == 2)
{
d = x * x * x * x * x * x * x * 8.0 + x * x * x * 60.0;
}
// 4. cos(x)
else if (Key_s == 3)
{
thrust::complex<float> c = thrust::complex<float>(x.real, x.imag);
d = {
cos(c).real(),
cos(c).imag()
};
}
return d;
}
__host__ __device__ float newton(ucomplex x0, float eps, int maxiter, int Key_d, int Key_s) {
ucomplex x = x0;
int iter = 0;
// 픽셀의 고유 iter값 계산
while (abs(f(x, Key_d, Key_s)) > eps && iter <= maxiter) {
iter++;
x = x - f(x, Key_d, Key_s) / df(x, Key_d, Key_s); // 점화식
}
return iter;
}
int main(int argc, char** argv) {
// opengl 초기화
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGB);
//glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGB);
// 윈도우 설정
glutInitWindowSize(width, width);
//glutInitWindowSize(700, 700);
glutInitWindowPosition(0, 0);
glutCreateWindow("Newton");
glClearColor(1.0, 1.0, 1.0, 1.0);
// 관측설정
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
glOrtho(-2, 2, 2, -2, 1, -1);
// 콜백함수 등록
glutDisplayFunc(Display);
glutTimerFunc(30, Timer, 0);
glutReshapeFunc(Reshape);
glutKeyboardFunc(Keyboard);
glutMouseFunc(Mouse);
glutMotionFunc(MouseMove);
glutMouseWheelFunc(MouseWheel);
// grid, block 만들기
const int block_width = 16;
Size = width * width * sizeof(float);
cudaMalloc((void**)&zD, Size);
cudaMalloc((void**)&DevImage, width * width * 3);
dimGrid = dim3((width - 1 / block_width) + 1, (width - 1 / block_width) + 1);
dimBlock = dim3(block_width, block_width);
p = (width / 1024);
glutMainLoop();
// 동적할당 해제
delete[] z;
cudaFree(DevImage);
cudaFree(zD);
}
void MathGPU() {
float scopenum = p * 15 * sin(scope);
float dx = Dx, dy = Dy;
cudaError_t cudaStatus = cudaSetDevice(0);
MathGPUKernel << <dimGrid, dimBlock >> > (zD, scopenum, dx, dy, key_d, key_s, DevImage);
cudaDeviceSynchronize();
cudaMemcpy(Image, DevImage, width * width * 3, cudaMemcpyDeviceToHost);
}
__global__ void MathGPUKernel(float* zD, float scopenum, float dx, float dy, int Key_d, int Key_s, unsigned char* DevImage) {
float xi = blockIdx.x * blockDim.x + threadIdx.x;
float yi = blockIdx.y * blockDim.y + threadIdx.y;
int maxiter = 255;
float eps = 0.0001;
// 뉴턴 점화식 계산
if (xi < width && yi < width) {
ucomplex xy = { (xi / width * 4 - 2) * scopenum + dx,(yi / width * 4 - 2) * scopenum + dy };
zD[(int)yi * width + (int)xi] = newton(xy, eps, maxiter, Key_d, Key_s);
// 색상 계산
float color = zD[(int)yi * width + (int)xi] / maxiter;
int offset = ((int)yi * width + (int)xi) * 3;
DevImage[offset] = color * 13 * 255;
DevImage[offset + 1] = color * 33 * 255;
DevImage[offset + 2] = color * 49 * 255;
//printf("%f\n", DevImage[offset]);
}
}
// width * width 픽셀에 색상을 지정
void Display() {
glClearColor(1, 1, 1, 1);
glClear(GL_COLOR_BUFFER_BIT);
__int64 _st = GetMicroSecond();
MathGPU();
//for (int i = 0; i < width * width * 3; i++)
// printf("%f\n", Image[i]);
glDrawPixels(width, width, GL_RGB, GL_UNSIGNED_BYTE, Image);
printf("Elapsed time: %I64d mico sec \n", GetMicroSecond() - _st);
glFinish();
}
void Timer(int id)
{
// Mode = true 일때 시간에따라 확대, 축소
if (Mode) {
if (TimeDir == true) {
scope -= 0.05;
if (scope <= 0.00001) {
scope = 0.00001;
TimeDir = false;
}
}
else {
scope += 0.05;
if (scope >= 1.57) {
scope = 1.57;
TimeDir = true;
}
}
}
glutPostRedisplay();
glutTimerFunc(30, Timer, 0); //30ms마다 Timer함수 호출
}
void Reshape(int w, int h)
{
glViewport(0, 0, w, h);
}
void Keyboard(unsigned char key, int x, int y)
{
if (key == 27)
exit(-1);
// 방정식 바꾸기
if (key == 's')
{
key_d = 0;
key_s = (key_s + 1) % 4;
}
// 차수 늘리기
if (key == 'd')
{
// 삼각함수는 차수 늘리지 않음
if (key_s != 3)
key_d++;
}
// 모드 바꾸기
if (key == 'f') {
Mode = !Mode;
}
glutPostRedisplay();
}
void Mouse(int button, int state, int x, int y)
{
// 마우스 버튼을 누른 경우,
if (state == GLUT_DOWN)
{
StartPt[0] = x;
StartPt[1] = y;
}
// 마우스 버튼을 땐 경우,
if (state == GLUT_UP)
StartPt[0] = StartPt[1] = 0;
}
void MouseMove(int x, int y) {
Dx += -(x - StartPt[0]) * 0.01; // x 이동값
Dy += -(StartPt[1] - y) * 0.01; // y 이동값
}
void MouseWheel(int button, int dir, int x, int y)
{
// Mode = false 일때
if (!Mode) {
if (dir > 0)
{
pre_scope = scope;
scope -= 0.01;
if (scope <= 0.00001)
scope = 0.00001;
}
else
{
pre_scope = scope;
scope += 0.01;
if (scope >= 1.57)
scope = 1.57;
}
}
glutPostRedisplay();
}
__int64 GetMicroSecond()
{
LARGE_INTEGER frequency;
LARGE_INTEGER now;
if (!QueryPerformanceFrequency(&frequency))
return (__int64)GetTickCount();
if (!QueryPerformanceCounter(&now))
return (__int64)GetTickCount();
return((now.QuadPart) / (frequency.QuadPart / 1000000));
}
*/-
CPU버전과 달리 2중 for문을 사용하지 않고 커널함수를 호출하여 width*width 만큼의 스레드를 생성하였다.
-
각 스레드별로 MathGpuKernel()함수 내에서 복소수 할당 -> newton점화식 계산 -> 색상 지정 까지의 과정이 한번에 진행된다.
-
PBO를 이용하여 구현하여 결과값을 CPU로 전달할 필요없이 GPU에서 바로 색상을 출력하여 성능개선을 하였다.
-
정점을 이용하여 출력 VS 픽셀을 이용하여 출력 과정을 비교해본 결과 픽셀을 이용해 출력하는 과정이 시간이 덜 걸리는 것을 확인 할 수 있었다. (주석 처리된 부분이 정점을 이용해 출력한 부분)
4. 결과 영상
5. 성능 비교
'수업 과제 & 실습 > 병렬프로그래밍' 카테고리의 다른 글
| GPU를 이용한 Julia Set 구현 (0) | 2020.07.04 |
|---|
'수업 과제 & 실습/병렬프로그래밍' Related Articles
more
