From 291ca0bd705edfee96d6514782b6183aebc606e7 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 15 May 2015 14:46:23 +0200 Subject: [PATCH 01/87] add antialiasing option --- hscgv_uebung_5/ApplicationWindow.cpp | 11 +++++------ hscgv_uebung_5/ApplicationWindow.h | 5 +++-- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/hscgv_uebung_5/ApplicationWindow.cpp b/hscgv_uebung_5/ApplicationWindow.cpp index fb104a9..bde899f 100644 --- a/hscgv_uebung_5/ApplicationWindow.cpp +++ b/hscgv_uebung_5/ApplicationWindow.cpp @@ -41,10 +41,10 @@ ApplicationWindow::ApplicationWindow() connect(ui.actionResetLight, SIGNAL(triggered()), this, SIGNAL(resetLight())); connect(ui.actionAnimate, SIGNAL(toggled(bool)), this, SLOT(animate(bool))); + connect(ui.actionAntialiasing, SIGNAL(toggled(bool)), SLOT(antialiasing(bool))); connect(ui.actionCPU, SIGNAL(toggled(bool)), SLOT(renderModeChanged())); connect(ui.actionGPU, SIGNAL(toggled(bool)), SLOT(renderModeChanged())); - connect(ui.actionAntialiasing, SIGNAL(toggled(bool)), SLOT(antialiasingChanged())); // ----- SIGNALS ----- @@ -156,6 +156,10 @@ void ApplicationWindow::animate(bool on) m_trigger->stop(); } +void ApplicationWindow::antialiasing(bool on) +{ + m_frame->m_antialiasing = on; +} int ApplicationWindow::getRenderMode() const { @@ -170,8 +174,3 @@ void ApplicationWindow::renderModeChanged() const emit renderMode(getRenderMode()); } -void ApplicationWindow::antialiasingChanged() const -{ - // TODO: change to appropriate antialiasing set method - emit renderMode(getRenderMode()); -} diff --git a/hscgv_uebung_5/ApplicationWindow.h b/hscgv_uebung_5/ApplicationWindow.h index 35a96d5..b21c6e5 100644 --- a/hscgv_uebung_5/ApplicationWindow.h +++ b/hscgv_uebung_5/ApplicationWindow.h @@ -96,14 +96,15 @@ class ApplicationWindow : public QMainWindow //! start and stop render timer void animate(bool on); + //! we want to have some nice output + void antialiasing(bool on); + //! deduce render mode from menu state int getRenderMode() const; //! make sure that render mode signal is triggerd void renderModeChanged() const; - //! we want to have some nice output - void antialiasingChanged() const; private: //! store status bar message From a0ad27ff59aab693cab6b937904e2e04fa6b83d0 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 15 May 2015 14:48:02 +0200 Subject: [PATCH 02/87] dynamic parameter setting via ctor --- hscgv_uebung_5/raytracer.cpp | 16 ++++++++-------- hscgv_uebung_5/raytracer.h | 2 +- 2 files changed, 9 insertions(+), 9 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cpp b/hscgv_uebung_5/raytracer.cpp index 452d22a..bc29c9a 100644 --- a/hscgv_uebung_5/raytracer.cpp +++ b/hscgv_uebung_5/raytracer.cpp @@ -43,16 +43,16 @@ Raytracer::Raytracer(bool antialiasing): m_filename(0), m_antialiasing(antialiasing) {} Raytracer::Raytracer(const char *filename, bool antialiasing): - m_isFileLoaded(true), m_filename(filename), m_antialiasing(antialiasing) { // parse the input file ReadScene(m_filename); + m_isFileLoaded = true; } void -Raytracer::start(float *renderedScene) { +Raytracer::start(float *renderedScene, int xRes, int yRes) { // setup viewport, its origin is bottom left // setup camera coordsys Vec3d eye_dir = (g_scene.view.lookat - g_scene.view.eyepoint).getNormalized(); @@ -64,15 +64,15 @@ Raytracer::start(float *renderedScene) { float width = height * g_scene.view.aspect; // compute delta steps in each direction - Vec3d deltaX = eye_right * (width / g_scene.picture.Xresolution); - Vec3d deltaY = eye_up * (height / g_scene.picture.Yresolution); + Vec3d deltaX = eye_right * (width / xRes); + Vec3d deltaY = eye_up * (height / yRes); // this should be bottom left - Vec3d bottomLeft = g_scene.view.eyepoint + eye_dir - deltaX*g_scene.picture.Xresolution/2 - deltaY*g_scene.picture.Yresolution/2; + Vec3d bottomLeft = g_scene.view.eyepoint + eye_dir - deltaX*xRes/2 - deltaY*yRes/2; // normal ray tracing: the color of the center of a pixel is computed - for (unsigned int sy=g_scene.picture.Yresolution ; sy > 0 ; --sy) { - for (unsigned int sx=0 ; sx < g_scene.picture.Xresolution ; ++sx) { + for (int sy=yRes ; sy > 0 ; --sy) { + for (int sx=0 ; sx < xRes ; ++sx) { // the center of the pixel we are looking at right now Vec3d point = bottomLeft + deltaX*sx + deltaY*sy + deltaX/2 + deltaY/2; @@ -110,7 +110,7 @@ Raytracer::start(float *renderedScene) { } } - int index = 3*((sy-1) * g_scene.picture.Xresolution + sx); + int index = 3*((sy-1) * xRes + sx); renderedScene[index + 0] = col[0]; renderedScene[index + 1] = col[1]; renderedScene[index + 2] = col[2]; diff --git a/hscgv_uebung_5/raytracer.h b/hscgv_uebung_5/raytracer.h index 73c61b3..9bed4e0 100644 --- a/hscgv_uebung_5/raytracer.h +++ b/hscgv_uebung_5/raytracer.h @@ -26,7 +26,7 @@ class Raytracer Raytracer(); Raytracer(bool antialiasing); Raytracer(const char* filename, bool antialiasing); - void start(float* renderedScene); + void start(float* renderedScene, int xRes, int yRes); bool m_isFileLoaded; private: const char *m_filename; From e67eb61ac0b8752be2a447b081a0332e42b2fc49 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 15 May 2015 14:50:28 +0200 Subject: [PATCH 03/87] render dynamically --- hscgv_uebung_5/GLFrame.cpp | 47 +++++++++++++++++++++----------------- hscgv_uebung_5/GLFrame.h | 9 +++++++- 2 files changed, 34 insertions(+), 22 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 38b1dea..5227c15 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -20,6 +20,7 @@ GLFrame::GLFrame(ApplicationWindow *parent) : QGLWidget(parent) +, m_antialiasing(false) , m_axesVisible(true) , m_modelVisible(true) , m_frameCounter(0) @@ -27,7 +28,7 @@ GLFrame::GLFrame(ApplicationWindow *parent) , m_renderMode(CPU) { // set minium size of rendering area - setMinimumSize(300,300); + setMinimumSize(150,150); // setup OpenGL buffers QGLFormat format; @@ -88,7 +89,7 @@ void GLFrame::loadTexture() glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER); // TODO adaptive size - glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, 600 , 600, 0, GL_RGB, GL_UNSIGNED_BYTE, (GLvoid*)m_data);CHECKGL; + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, m_width, m_height, 0, GL_RGB, GL_FLOAT, (GLvoid*)m_data);CHECKGL; glBindTexture(GL_TEXTURE_2D, 0); @@ -149,17 +150,7 @@ void GLFrame::paintGL() if(m_axesVisible) drawCoordSys(); - switch(m_renderMode) - { - case CPU: - drawScene(m_renderMode); - break; - case GPU: - drawScene(m_renderMode); - break; - } - - loadTexture(); + renderScene(m_renderMode); drawFullScreenQuad(); // draw transparent light source last @@ -426,6 +417,12 @@ void GLFrame::setAxesVisibility(bool on) updateGL(); } +void GLFrame::setAntialiasing(bool on) +{ + m_antialiasing = on; + updateGL(); +} + // reset camera to default values void GLFrame::resetLight() { @@ -519,30 +516,38 @@ void GLFrame::drawFullScreenQuad() } -// draw model -void GLFrame::drawScene(RenderMode mode) +// render scene and save it into a texture +void GLFrame::renderScene(RenderMode mode) { + // only render the scene if there is some intialized raytracer if(!m_raytracer || !m_raytracer->m_isFileLoaded) return; - // Draw Model/Scene + // draw scene switch(mode) { case CPU: - m_data = (float*)malloc(sizeof(float) * 3 * m_height * m_width); - m_raytracer->start(m_data); - free(m_data); + m_data = (float*)malloc(sizeof(float) * 3 * m_width * m_height); + m_raytracer->start(m_data, m_width, m_height); + loadTexture(); +// updateGL(); break; case GPU: break; } } +void GLFrame::loadViewFromData() +{ + + +} + void GLFrame::loadScene(const QString &filename) { delete m_raytracer; - - m_raytracer = new Raytracer(filename.toStdString().c_str(),false); + m_raytracer = new Raytracer(filename.toStdString().c_str(), m_antialiasing); + loadViewFromData(); updateGL(); } diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index 8337dc4..394e421 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -68,6 +68,9 @@ class GLFrame : public QGLWidget //! accessor for frame counter int frameCounter() const; + //! antialiasing + bool m_antialiasing; + // (Some Slots and Signals) signals: //! show some information to the user @@ -78,12 +81,16 @@ public slots: void setRenderMode(int mode); //! set visibility of coordinate axes void setAxesVisibility(bool on); + //! set antialising + void setAntialiasing(bool on); //! reset camera to default values void resetCam(); //! reset light to default values void resetLight(); //! read scene file void loadScene(const QString &filename); + //! load the viewport, camera, position etc. into this frame + void loadViewFromData(); protected: // ---- OGLWidget Basic Methods ---- @@ -118,7 +125,7 @@ public slots: void drawFullScreenQuad(); //! draw scene - void drawScene(RenderMode mode); + void renderScene(RenderMode mode); //! load current raytracing scene into texture void loadTexture(); From 7ac16a015ff478c61a3113579af44e14d3558ea5 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 15 May 2015 15:10:33 +0200 Subject: [PATCH 04/87] handle mac and unix build at the same time --- hscgv_uebung_5/GLFrame.h | 5 +++++ hscgv_uebung_5/src.pro | 2 +- 2 files changed, 6 insertions(+), 1 deletion(-) diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index 394e421..84db8a0 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -10,8 +10,13 @@ #ifndef GLFRAME_H #define GLFRAME_H +#ifdef __APPLE__ +#include +#include +#elif #include #include +#endif #include #include #include diff --git a/hscgv_uebung_5/src.pro b/hscgv_uebung_5/src.pro index c9ba74a..0c88566 100644 --- a/hscgv_uebung_5/src.pro +++ b/hscgv_uebung_5/src.pro @@ -27,7 +27,7 @@ QMAKE_CXXFLAGS += -W -Wall FORMS += \ ApplicationWindow.ui -unix:LIBS *= -lGLEW -lglut -lGLU +unix:!macx:LIBS *= -lGLEW -lglut -lGLU macx { LIBS *= -lGLEW -L/usr/local/opt/glew/lib From 2e87ce6e2f18bf11b66c26f07400c0f66e73c2b3 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 15 May 2015 15:59:07 +0200 Subject: [PATCH 05/87] nasty bug removed --- hscgv_uebung_5/raytracer.cpp | 8 ++++---- hscgv_uebung_5/raytracer.h | 1 + 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cpp b/hscgv_uebung_5/raytracer.cpp index bc29c9a..647018c 100644 --- a/hscgv_uebung_5/raytracer.cpp +++ b/hscgv_uebung_5/raytracer.cpp @@ -50,6 +50,10 @@ Raytracer::Raytracer(const char *filename, bool antialiasing): ReadScene(m_filename); m_isFileLoaded = true; } +Raytracer::~Raytracer() { + // clean up + cleanUp(); +} void Raytracer::start(float *renderedScene, int xRes, int yRes) { @@ -116,8 +120,4 @@ Raytracer::start(float *renderedScene, int xRes, int yRes) { renderedScene[index + 2] = col[2]; } // foreach x } // foreach y - - // clean up - cleanUp(); - } diff --git a/hscgv_uebung_5/raytracer.h b/hscgv_uebung_5/raytracer.h index 9bed4e0..d1859b0 100644 --- a/hscgv_uebung_5/raytracer.h +++ b/hscgv_uebung_5/raytracer.h @@ -26,6 +26,7 @@ class Raytracer Raytracer(); Raytracer(bool antialiasing); Raytracer(const char* filename, bool antialiasing); + ~Raytracer(); void start(float* renderedScene, int xRes, int yRes); bool m_isFileLoaded; private: From 5a9c74657230bf7146d4da751327f45b4a188816 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 15 May 2015 16:00:05 +0200 Subject: [PATCH 06/87] add raytracingNeeded flag --- hscgv_uebung_5/GLFrame.cpp | 17 +++++++++++++++-- hscgv_uebung_5/GLFrame.h | 1 + 2 files changed, 16 insertions(+), 2 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 5227c15..4be920c 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -25,6 +25,7 @@ GLFrame::GLFrame(ApplicationWindow *parent) , m_modelVisible(true) , m_frameCounter(0) , m_raytracer(NULL) +, m_raytracingNeeded(false) , m_renderMode(CPU) { // set minium size of rendering area @@ -104,6 +105,7 @@ void GLFrame::resizeGL(int w, int h) m_height = h; glViewport(0, 0, (GLint)w, (GLint)h); + m_raytracingNeeded = true; updateGL(); } @@ -197,6 +199,7 @@ void GLFrame::mouseMoveEvent(QMouseEvent *m) m_lastMouseY = mouseY; // force openGL to redraw the changes + m_raytracingNeeded = true; updateGL(); } @@ -208,6 +211,7 @@ void GLFrame::wheelEvent(QWheelEvent *ev) else adjustCam(false, false, true, 0, -0.001*ev->delta()); + m_raytracingNeeded = true; updateGL(); } @@ -357,6 +361,8 @@ void GLFrame::adjustLight(bool leftButton, bool middleButton, bool rightButton, matMult(inc, invCam, inc); matMult(m_lightRot, inc, m_lightRot); } + // TODO + m_raytracingNeeded = true; } // interpret mouse movement in order to move the camera or adjust @@ -402,12 +408,15 @@ void GLFrame::adjustCam(bool leftButton, bool middleButton, bool rightButton, if(m_fieldOfView > 160.0) m_fieldOfView = 160.0; } + // TODO + m_raytracingNeeded = true; } // notify when the drawMode should be changed void GLFrame::setRenderMode(int mode) { m_renderMode = (RenderMode)mode; + m_raytracingNeeded = true; updateGL(); } @@ -420,6 +429,7 @@ void GLFrame::setAxesVisibility(bool on) void GLFrame::setAntialiasing(bool on) { m_antialiasing = on; + m_raytracingNeeded = true; updateGL(); } @@ -430,6 +440,7 @@ void GLFrame::resetLight() matIdent(m_lightRot); + m_raytracingNeeded = true; updateGL(); } @@ -441,6 +452,7 @@ void GLFrame::resetCam() matIdent(m_camRot); m_fieldOfView = 15.0; + m_raytracingNeeded = true; updateGL(); } @@ -520,7 +532,7 @@ void GLFrame::drawFullScreenQuad() void GLFrame::renderScene(RenderMode mode) { // only render the scene if there is some intialized raytracer - if(!m_raytracer || !m_raytracer->m_isFileLoaded) + if(!m_raytracingNeeded || !m_raytracer || !m_raytracer->m_isFileLoaded) return; // draw scene @@ -529,8 +541,8 @@ void GLFrame::renderScene(RenderMode mode) case CPU: m_data = (float*)malloc(sizeof(float) * 3 * m_width * m_height); m_raytracer->start(m_data, m_width, m_height); + m_raytracingNeeded = false; loadTexture(); -// updateGL(); break; case GPU: break; @@ -547,6 +559,7 @@ void GLFrame::loadScene(const QString &filename) { delete m_raytracer; m_raytracer = new Raytracer(filename.toStdString().c_str(), m_antialiasing); + m_raytracingNeeded = true; loadViewFromData(); updateGL(); } diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index 84db8a0..da5fed8 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -187,6 +187,7 @@ public slots: //! currently loaded scene Raytracer *m_raytracer; + bool m_raytracingNeeded; float* m_data; //! how our model is to be drawn From a06e8478779fc9c1d246c73251a332f96b2c1e67 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 15 May 2015 16:00:21 +0200 Subject: [PATCH 07/87] free user data and update todos --- hscgv_uebung_5/GLFrame.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 4be920c..aa99e22 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -89,9 +89,12 @@ void GLFrame::loadTexture() glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER); - // TODO adaptive size + // move data into a texture glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, m_width, m_height, 0, GL_RGB, GL_FLOAT, (GLvoid*)m_data);CHECKGL; + // free data since its on the texture + free(m_data); + glBindTexture(GL_TEXTURE_2D, 0); } @@ -549,6 +552,7 @@ void GLFrame::renderScene(RenderMode mode) } } +// TODO load initial position of camera, viewport etc. void GLFrame::loadViewFromData() { From 42431b118866ed0c1412e8407f84806980bde589 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 15 May 2015 17:14:32 +0200 Subject: [PATCH 08/87] parallelize CPU code --- hscgv_uebung_5/raytracer.cpp | 1 + hscgv_uebung_5/src.pro | 3 ++- 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/hscgv_uebung_5/raytracer.cpp b/hscgv_uebung_5/raytracer.cpp index 647018c..5bf2907 100644 --- a/hscgv_uebung_5/raytracer.cpp +++ b/hscgv_uebung_5/raytracer.cpp @@ -75,6 +75,7 @@ Raytracer::start(float *renderedScene, int xRes, int yRes) { Vec3d bottomLeft = g_scene.view.eyepoint + eye_dir - deltaX*xRes/2 - deltaY*yRes/2; // normal ray tracing: the color of the center of a pixel is computed + #pragma omp parallel for schedule(dynamic) collapse(2) for (int sy=yRes ; sy > 0 ; --sy) { for (int sx=0 ; sx < xRes ; ++sx) { // the center of the pixel we are looking at right now diff --git a/hscgv_uebung_5/src.pro b/hscgv_uebung_5/src.pro index 0c88566..2862f8d 100644 --- a/hscgv_uebung_5/src.pro +++ b/hscgv_uebung_5/src.pro @@ -22,7 +22,8 @@ DEFINES *= TRACE VERBOSE CONFIG += debug CONFIG -= release -QMAKE_CXXFLAGS += -W -Wall +QMAKE_CXXFLAGS += -W -Wall -fopenmp +#QMAKE_LFLAGS += -fopenmp FORMS += \ ApplicationWindow.ui From 2a8c2862bf11d287a753464ff0855b296d325040 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Sun, 17 May 2015 12:51:30 +0200 Subject: [PATCH 09/87] use scrolling to move in and out of the scene --- hscgv_uebung_5/GLFrame.cpp | 3 +++ hscgv_uebung_5/GLFrame.h | 1 + 2 files changed, 4 insertions(+) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index aa99e22..f6192ab 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -403,6 +403,9 @@ void GLFrame::adjustCam(bool leftButton, bool middleButton, bool rightButton, } else if(rightButton) { + // move eyepoint into scrolling direction + g_scene.view.eyepoint += Vec3d(.0,.0,dy*1000); + // prevent negative fieldOfView if(dy < -0.5) dy = 0.5; diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index da5fed8..33627db 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -21,6 +21,7 @@ #include #include #include "raytracer.h" +#include "param.h" Q_DECLARE_METATYPE(unsigned char *) From bce6484f4dee7d6288d37ea146794d545a615513 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Sun, 17 May 2015 12:51:54 +0200 Subject: [PATCH 10/87] remove initial data loading - not needed --- hscgv_uebung_5/GLFrame.cpp | 7 ------- hscgv_uebung_5/GLFrame.h | 2 -- 2 files changed, 9 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index f6192ab..2c4e17a 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -555,19 +555,12 @@ void GLFrame::renderScene(RenderMode mode) } } -// TODO load initial position of camera, viewport etc. -void GLFrame::loadViewFromData() -{ - - -} void GLFrame::loadScene(const QString &filename) { delete m_raytracer; m_raytracer = new Raytracer(filename.toStdString().c_str(), m_antialiasing); m_raytracingNeeded = true; - loadViewFromData(); updateGL(); } diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index 33627db..84d33b5 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -95,8 +95,6 @@ public slots: void resetLight(); //! read scene file void loadScene(const QString &filename); - //! load the viewport, camera, position etc. into this frame - void loadViewFromData(); protected: // ---- OGLWidget Basic Methods ---- From be4caef186190c54a047fa30df2dfa4b430e4936 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Sun, 17 May 2015 13:15:34 +0200 Subject: [PATCH 11/87] remove action menu items --- hscgv_uebung_5/ApplicationWindow.ui | 3 --- 1 file changed, 3 deletions(-) diff --git a/hscgv_uebung_5/ApplicationWindow.ui b/hscgv_uebung_5/ApplicationWindow.ui index 87d2f0d..d67c5a3 100644 --- a/hscgv_uebung_5/ApplicationWindow.ui +++ b/hscgv_uebung_5/ApplicationWindow.ui @@ -36,9 +36,6 @@ View - - - From b886a7f54e5752b7a8e222b3fc6a108329223aa9 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Sun, 17 May 2015 13:15:53 +0200 Subject: [PATCH 12/87] rotate via dragging the scene --- hscgv_uebung_5/GLFrame.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 2c4e17a..8b156e9 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -394,6 +394,12 @@ void GLFrame::adjustCam(bool leftButton, bool middleButton, bool rightButton, // multiply viewing transformation with incremental transformation // to become the new viewing transformation matMult(m_camRot, inc, m_camRot); + Vec3d eyeg = Vec3d(g_scene.view.eyepoint); + GLfloat eye[4] = {eyeg[0],eyeg[1],eyeg[2],.0}; + matMult(eye, inc, eye); + for (int i=0;i<3;i++) + g_scene.view.eyepoint[i] = eye[i]; + } } else if(middleButton) From 4f8da96e2d315f1fe07a4e83576a55587da4b42c Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Sun, 17 May 2015 13:24:44 +0200 Subject: [PATCH 13/87] prevent negative fov --- hscgv_uebung_5/GLFrame.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 8b156e9..5c6f6f9 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -364,7 +364,6 @@ void GLFrame::adjustLight(bool leftButton, bool middleButton, bool rightButton, matMult(inc, invCam, inc); matMult(m_lightRot, inc, m_lightRot); } - // TODO m_raytracingNeeded = true; } @@ -409,13 +408,14 @@ void GLFrame::adjustCam(bool leftButton, bool middleButton, bool rightButton, } else if(rightButton) { - // move eyepoint into scrolling direction - g_scene.view.eyepoint += Vec3d(.0,.0,dy*1000); // prevent negative fieldOfView if(dy < -0.5) dy = 0.5; + // move eyepoint into scrolling direction + g_scene.view.eyepoint += Vec3d(.0,.0,dy*1000); + m_fieldOfView *= 1.0+dy; if(m_fieldOfView > 160.0) m_fieldOfView = 160.0; From 0b092f3407553f6abfdc2239d6c5635a2501e80f Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Sun, 17 May 2015 13:24:56 +0200 Subject: [PATCH 14/87] smoother moving in scene --- hscgv_uebung_5/GLFrame.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 5c6f6f9..4db43ed 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -388,7 +388,7 @@ void GLFrame::adjustCam(bool leftButton, bool middleButton, bool rightButton, { float a[4] = {dy/d, -dx/d, 0.0, 0.0}; // rotation axis GLfloat inc[16]; - matRot(inc, a, d*10); + matRot(inc, a, d*5); // multiply viewing transformation with incremental transformation // to become the new viewing transformation From 34a15d29552f39f70ea589684177aff9c71b7a7e Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Sun, 17 May 2015 13:25:18 +0200 Subject: [PATCH 15/87] camera rotation is not needed --- hscgv_uebung_5/GLFrame.cpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 4db43ed..5b11470 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -390,9 +390,7 @@ void GLFrame::adjustCam(bool leftButton, bool middleButton, bool rightButton, GLfloat inc[16]; matRot(inc, a, d*5); - // multiply viewing transformation with incremental transformation - // to become the new viewing transformation - matMult(m_camRot, inc, m_camRot); + // multiply transformation matrix with eyepoint to get new perspective Vec3d eyeg = Vec3d(g_scene.view.eyepoint); GLfloat eye[4] = {eyeg[0],eyeg[1],eyeg[2],.0}; matMult(eye, inc, eye); @@ -420,7 +418,6 @@ void GLFrame::adjustCam(bool leftButton, bool middleButton, bool rightButton, if(m_fieldOfView > 160.0) m_fieldOfView = 160.0; } - // TODO m_raytracingNeeded = true; } From b473c873809311de365a4f8e6c3440572c67cd47 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Sun, 17 May 2015 13:29:07 +0200 Subject: [PATCH 16/87] remove mode, since it's a member --- hscgv_uebung_5/GLFrame.cpp | 6 +++--- hscgv_uebung_5/GLFrame.h | 2 +- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 5b11470..8c334de 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -155,7 +155,7 @@ void GLFrame::paintGL() if(m_axesVisible) drawCoordSys(); - renderScene(m_renderMode); + renderScene(); drawFullScreenQuad(); // draw transparent light source last @@ -538,14 +538,14 @@ void GLFrame::drawFullScreenQuad() } // render scene and save it into a texture -void GLFrame::renderScene(RenderMode mode) +void GLFrame::renderScene() { // only render the scene if there is some intialized raytracer if(!m_raytracingNeeded || !m_raytracer || !m_raytracer->m_isFileLoaded) return; // draw scene - switch(mode) + switch(m_renderMode) { case CPU: m_data = (float*)malloc(sizeof(float) * 3 * m_width * m_height); diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index 84d33b5..b1d53fe 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -129,7 +129,7 @@ public slots: void drawFullScreenQuad(); //! draw scene - void renderScene(RenderMode mode); + void renderScene(); //! load current raytracing scene into texture void loadTexture(); From 8d49d4e2183b2954dc5d4b13c707040e30b42c1c Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Sun, 17 May 2015 16:01:55 +0200 Subject: [PATCH 17/87] prepare for cuda implementation --- hscgv_uebung_5/GLFrame.cpp | 8 + hscgv_uebung_5/cuda.pri | 34 ++++ .../{geoobject.cpp => geoobject.cu} | 24 +-- .../{geoquadric.cpp => geoquadric.cu} | 10 +- hscgv_uebung_5/ray.cu | 167 ++++++++++++++++++ hscgv_uebung_5/raytracer.cu | 69 ++++++++ hscgv_uebung_5/src.pro | 15 +- 7 files changed, 309 insertions(+), 18 deletions(-) create mode 100644 hscgv_uebung_5/cuda.pri rename hscgv_uebung_5/{geoobject.cpp => geoobject.cu} (81%) rename hscgv_uebung_5/{geoquadric.cpp => geoquadric.cu} (89%) create mode 100644 hscgv_uebung_5/ray.cu create mode 100644 hscgv_uebung_5/raytracer.cu diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 8c334de..2e770b1 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -425,6 +425,14 @@ void GLFrame::adjustCam(bool leftButton, bool middleButton, bool rightButton, void GLFrame::setRenderMode(int mode) { m_renderMode = (RenderMode)mode; + + if(mode == CPU) { + + } + else if (mode == GPU) { + // Todo: setup cuda (maybe earlier) and push all data to GPU + } + m_raytracingNeeded = true; updateGL(); } diff --git a/hscgv_uebung_5/cuda.pri b/hscgv_uebung_5/cuda.pri new file mode 100644 index 0000000..0b653ac --- /dev/null +++ b/hscgv_uebung_5/cuda.pri @@ -0,0 +1,34 @@ +######################################################################## +# CUDA +# see: http://forums.nvidia.com/index.php?showtopic=29539 +######################################################################## +win32 { + INCLUDEPATH += $(CUDA_INC_DIR) + QMAKE_LIBDIR += $(CUDA_LIB_DIR) + LIBS += -lcudart + + cuda.output = $$OBJECTS_DIR/${QMAKE_FILE_BASE}_cuda.obj + cuda.commands = $(CUDA_BIN_DIR)/nvcc.exe -c -Xcompiler $$join(QMAKE_CXXFLAGS,",") $$join(INCLUDEPATH,'" -I "','-I "','"') ${QMAKE_FILE_NAME} -o ${QMAKE_FILE_OUT} +} + +unix { +# auto-detect CUDA path + CUDA_DIR = $$system(which nvcc | sed 's,/bin/nvcc$,,') + CUDA_DIR = /usr/local/cuda-5.5 + INCLUDEPATH += $$CUDA_DIR/include + QMAKE_LIBDIR += $$CUDA_DIR/lib64 + LIBS += -lcudart + NVCCFLAGS="-use_fast_math --ptxas-options=-v" + NVCCFLAGS+="-Xptxas -fastimul -arch=sm_20" # 24 bit integer multiplication should be sufficient + debug:NVCCFLAGS+="-g" + release:NVCCFLAGS+="-O3" + macx:NVCCFLAGS+="-m64" + FLAGS=$$join(QMAKE_CXXFLAGS,",") $$join(INCLUDEPATH,'" -I "','-I "','"') $$join(DEFINES,'" -D"','-D"','"') + cuda.output = ${OBJECTS_DIR}${QMAKE_FILE_BASE}_cuda.o + cuda.commands = $$CUDA_DIR/bin/nvcc -c $$NVCCFLAGS -Xcompiler $$FLAGS ${QMAKE_FILE_NAME} -o ${QMAKE_FILE_OUT} + + cuda.depend_command = $$CUDA_DIR/bin/nvcc -M $$NVCCFLAGS -Xcompiler $$FLAGS ${QMAKE_FILE_NAME} | sed "'s,^.*: ,,'" | tr -d '\\\\\\n' +} +cuda.input = CUDA_SOURCES +QMAKE_EXTRA_COMPILERS += cuda +######################################################################## diff --git a/hscgv_uebung_5/geoobject.cpp b/hscgv_uebung_5/geoobject.cu similarity index 81% rename from hscgv_uebung_5/geoobject.cpp rename to hscgv_uebung_5/geoobject.cu index 08c93dc..dbfe970 100644 --- a/hscgv_uebung_5/geoobject.cpp +++ b/hscgv_uebung_5/geoobject.cu @@ -8,30 +8,31 @@ /* member access functions for GeoObjectProperties */ -Color GeoObjectProperties::ambient() const +Color __host__ __device__ +GeoObjectProperties::ambient() const { return m_ambient; } -Vec3d +Vec3d __host__ __device__ GeoObjectProperties::reflectance() const { return m_reflectance; } -double +double __host__ __device__ GeoObjectProperties::specular() const { return m_specular; } -int +int __host__ __device__ GeoObjectProperties::specularExp() const { return m_specularExp; } -double +double __host__ __device__ GeoObjectProperties::mirror() const { return m_mirror; @@ -39,7 +40,8 @@ GeoObjectProperties::mirror() const /* access functions for the properties of a GeoObject */ -Color GeoObject::ambient() const +Color __host__ __device__ +GeoObject::ambient() const { if (m_properties) return m_properties->ambient(); @@ -48,7 +50,7 @@ Color GeoObject::ambient() const return Vec3d(0.0); } -Vec3d +Vec3d __host__ __device__ GeoObject::reflectance() const { if (m_properties) @@ -58,7 +60,7 @@ GeoObject::reflectance() const return Vec3d(0.0); } -double +double __host__ __device__ GeoObject::specular() const { if (m_properties) @@ -68,7 +70,7 @@ GeoObject::specular() const return 0.0; } -double +double __host__ __device__ GeoObject::specularExp() const { if (m_properties) @@ -78,7 +80,7 @@ GeoObject::specularExp() const return 0.0; } -double +double __host__ __device__ GeoObject::mirror() const { if (m_properties) @@ -88,7 +90,7 @@ GeoObject::mirror() const return 0.0; } -void +void __host__ __device__ GeoObject::setProperties(GeoObjectProperties *p) { m_properties = p; diff --git a/hscgv_uebung_5/geoquadric.cpp b/hscgv_uebung_5/geoquadric.cu similarity index 89% rename from hscgv_uebung_5/geoquadric.cpp rename to hscgv_uebung_5/geoquadric.cu index 47df36a..835022b 100644 --- a/hscgv_uebung_5/geoquadric.cpp +++ b/hscgv_uebung_5/geoquadric.cu @@ -10,7 +10,7 @@ // Description: // Constructor. -GeoQuadric::GeoQuadric() + __host__ __device__ GeoQuadric::GeoQuadric() : m_a(0), m_b(0), m_c(0) , m_d(0), m_e(0), m_f(0) , m_g(0), m_h(0), m_j(0) @@ -21,7 +21,7 @@ GeoQuadric::GeoQuadric() // Description: // explicit parametrization -GeoQuadric::GeoQuadric(double na, double nb, double nc, double nd, double ne, + __host__ __device__ GeoQuadric::GeoQuadric(double na, double nb, double nc, double nd, double ne, double nf, double ng, double nh, double nj, double nk) : m_a(na), m_b(nb), m_c(nc) , m_d(nd), m_e(ne), m_f(nf) @@ -33,14 +33,14 @@ GeoQuadric::GeoQuadric(double na, double nb, double nc, double nd, double ne, // Description: // Destructor. -GeoQuadric::~GeoQuadric() + __host__ __device__ GeoQuadric::~GeoQuadric() { } // Description: // Return the normal vector at point v of this surface -Vec3d +Vec3d __host__ __device__ GeoQuadric::getNormal(const Vec3d &v) const { Vec3d tmpvec( (v | Vec3d(2*m_a,m_b,m_c)) + m_d, @@ -52,7 +52,7 @@ GeoQuadric::getNormal(const Vec3d &v) const // Description: // Compute intersection point on this surface // return distance to nearest intersection found or -1 -double +double __host__ __device__ GeoQuadric::intersect(const Ray &ray) const { double t = -1.0, acoef, bcoef, ccoef, root, disc; diff --git a/hscgv_uebung_5/ray.cu b/hscgv_uebung_5/ray.cu new file mode 100644 index 0000000..f3d396a --- /dev/null +++ b/hscgv_uebung_5/ray.cu @@ -0,0 +1,167 @@ +/* ******** Programmierpraktikum Computergrafik (CGP) ************** + * Aufgabe 1 - "Lichtblick" + * Created by Peter Kipfer + * Changed by Martin Aumueller + */ + +#include +#include +#include +#include "ray.h" +#include "param.h" +#include "types.h" + +// Description: +// Constructor. +Ray::Ray( ) +: m_origin() +, m_direction() +, m_depth(0) +, m_objList(NULL) +, m_lightList(NULL) +{ + ERR("don't use uninitialized rays !"); +} + +// Description: +// Copy-Constructor. +Ray::Ray( const Ray &r ) +{ + // copy it ! initialization is not enough ! + m_origin = r.m_origin; + m_direction = r.m_direction; + m_depth = r.m_depth; + m_objList = r.m_objList; + m_lightList = r.m_lightList; +} + +// Description: +// Constructor with explicit parameters +Ray::Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll) +{ + // copy it ! initialization is not enough ! + m_origin = o; + m_direction = d; + m_depth = i; + m_objList = &ol; + m_lightList = ≪ +} + +// Description: +// Destructor. +Ray::~Ray( ) +{ +} + +// Description: +// Determine color of this ray by tracing through the scene +const Color __device__ +Ray::shade() const +{ + for (int i=0; i<5; i++) { + GeoObject *closest = NULL; + double tMin = DBL_MAX; + + // find closest object that intersects + for (std::vector::iterator iter = m_objList->begin(); + iter != m_objList->end(); + iter++) { + + double t = (*iter)->intersect(*this); + if (0.0 < t && t < tMin) { + tMin = t; + closest = (*iter); + } + } + + // no object hit -> ray goes to infinity + if (closest == NULL) { + if (m_depth == 0) { + return g_scene.picture.background; // background color + } + else { + return Color(0.0); // black + } + } + else { + // reflection + Vec3d intersectionPosition(m_origin + (m_direction * tMin)); + Vec3d normal(closest->getNormal(intersectionPosition)); + Ray reflectedRay(intersectionPosition, + m_direction.getReflectedAt(normal).getNormalized(), + m_depth+1,*m_objList,*m_lightList); + Color currentColor(0.0); + + // calculate lighting + for (std::vector::iterator iter = m_lightList->begin(); + iter != m_lightList->end(); + iter++) { + + // where is the lightsource ? + Ray rayoflight(intersectionPosition, (*iter)->direction(), 0, *m_objList, *m_lightList); + bool something_intersected = false; + + // where are the objects ? + for (std::vector::iterator iter2 = m_objList->begin(); + iter2 != m_objList->end(); + iter2++) { + + double t = (*iter2)->intersect(rayoflight); + if (t > 0.0) { + something_intersected = true; + break; + } + + } // for all obj + + // is it visible ? + if (! something_intersected) + currentColor += shadedColor((*iter), reflectedRay, normal, closest); + + } // for all lights + + // could be right... + currentColor *= closest->mirror(); + } + } + return Color(currentColor); +} + +// Description: +// Determine color contribution of a lightsource +const Color __device__ __host__ +Ray::shadedColor(LightObject *light, const Ray &reflectedRay, const Vec3d &normal, GeoObject *obj) const +{ + double ldot = light->direction() | normal; + Color reflectedColor = Color(0.0); + + // lambertian reflection model + if (ldot > 0.0) + reflectedColor += obj->reflectance() * (light->color() * ldot); + + // updated with ambient lightning as in: + // [GENERALISED AMBIENT REFLECTION MODELS FOR LAMBERTIAN AND PHONG SURFACES, Xiaozheng Zhang and Yongsheng Gao] + reflectedColor += obj->ambient() * g_scene.ambience; + + // specular part + double spec = reflectedRay.direction() | light->direction(); + if (spec > 0.0) { + spec = obj->specular() * pow(spec, obj->specularExp()); + reflectedColor += light->color() * spec; + } + + return Color(reflectedColor); +} + +Vec3d __device__ __host__ +Ray::origin() const +{ + return Vec3d(m_origin); +} + +Vec3d __device__ __host__ +Ray::direction() const +{ + return Vec3d(m_direction); +} + diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu new file mode 100644 index 0000000..ad9f0f2 --- /dev/null +++ b/hscgv_uebung_5/raytracer.cu @@ -0,0 +1,69 @@ +#include "raytracer.h" + +void +Raytracer::start(float *renderedScene, int xRes, int yRes) { + // setup viewport, its origin is bottom left + // setup camera coordsys + Vec3d eye_dir = (g_scene.view.lookat - g_scene.view.eyepoint).getNormalized(); + Vec3d eye_right = (eye_dir^g_scene.view.up).getNormalized(); + Vec3d eye_up = eye_dir^eye_right*-1; + + // calculatehe dimensions of the viewport using the scene's camera + float height = 2 * tan(M_PI/180 * .5 * g_scene.view.fovy); + float width = height * g_scene.view.aspect; + + // compute delta steps in each direction + Vec3d deltaX = eye_right * (width / xRes); + Vec3d deltaY = eye_up * (height / yRes); + + // this should be bottom left + Vec3d bottomLeft = g_scene.view.eyepoint + eye_dir - deltaX*xRes/2 - deltaY*yRes/2; + + // normal ray tracing: the color of the center of a pixel is computed + for (int sy=yRes ; sy > 0 ; --sy) { + for (int sx=0 ; sx < xRes ; ++sx) { + // the center of the pixel we are looking at right now + Vec3d point = bottomLeft + deltaX*sx + deltaY*sy + deltaX/2 + deltaY/2; + + // the direction of our look + Vec3d dir = point - g_scene.view.eyepoint; + + // create ray from view.eyepoint to view.lookat + Ray theRay(g_scene.view.eyepoint,dir.getNormalized(),0,g_objectList,g_lightList); + + // compute the color + Color col = theRay.shade(); + + // in case we are using antialiasing, calculate the color of this pixel by averaging + if (m_antialiasing) { + // scale the midpoint color since we are going to use 5 points to average our color + col *= 0.2; + + // besides taking shooting a ray through the midpoint 'o', we calculate + // the pixels color by shooting 4 more rays through the points 'x' + // and averaging their values + // ------- + // | x x | + // | o | + // | x x | + // -------- + for(float dx = -1/4.; dx <= 1/4.; dx+=1/2.) { + for(float dy = -1/4.; dy <= 1/4.; dy+=1/2.) { + Vec3d superSamplePoint = point + deltaX*dx + deltaY*dy; + Vec3d superSampleDir = superSamplePoint - g_scene.view.eyepoint; + + // create ray from view.eyepoint to view.lookat + Ray theRay(g_scene.view.eyepoint,superSampleDir.getNormalized(),0,g_objectList,g_lightList); + col += theRay.shade()*0.2;//color, recursive_ray_trace(eye, ray, 0)); + } + } + } + + int index = 3*((sy-1) * xRes + sx); + renderedScene[index + 0] = col[0]; + renderedScene[index + 1] = col[1]; + renderedScene[index + 2] = col[2]; + } // foreach x + } // foreach y +} + diff --git a/hscgv_uebung_5/src.pro b/hscgv_uebung_5/src.pro index 2862f8d..9cdeebd 100644 --- a/hscgv_uebung_5/src.pro +++ b/hscgv_uebung_5/src.pro @@ -6,13 +6,18 @@ HEADERS= geoobject.h geoquadric.h lightobject.h ray.h types.h vector.h pa ApplicationWindow.h \ GLFrame.h \ raytracer.h -SOURCES= geoobject.cpp geoquadric.cpp lightobject.cpp main.cpp ray.cpp \ +SOURCES= lightobject.cpp main.cpp ray.cpp \ geopolygon.cpp \ ApplicationWindow.cpp \ GLFrame.cpp \ raytracer.cpp win32:SOURCES*= xgetopt.cpp YACCSOURCES= input.y +CUDA_SOURCES += \ + raytracer.cu \ + ray.cu \ + geoobject.cu \ + geoquadric.cu QMAKE_YACC = bison QMAKE_YACCFLAGS = -y -d @@ -28,7 +33,10 @@ QMAKE_CXXFLAGS += -W -Wall -fopenmp FORMS += \ ApplicationWindow.ui -unix:!macx:LIBS *= -lGLEW -lglut -lGLU +unix:!macx:{ +LIBS *= -lGLEW -lglut -lGLU +include(cuda.pri) +} macx { LIBS *= -lGLEW -L/usr/local/opt/glew/lib @@ -42,3 +50,6 @@ INCLUDEPATH *= $$GLEWDIR/include DEFINES *= GLEW_STATIC LIBS *= -L$$GLEWDIR/lib -lglew32s } + +OTHER_FILES += + From d5fe0638ec2de2eaa984e6aedba8b48f896a964f Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 18 May 2015 18:05:47 +0200 Subject: [PATCH 18/87] make it cuda ready - inline all device functions and move them into .inl files - mark necessary functions as host/device - setup only one class as cuda compilable file --- hscgv_uebung_5/geoobject.h | 27 +-- .../{geoobject.cu => geoobject.inl} | 42 ++--- hscgv_uebung_5/geoquadric.h | 3 +- .../{geoquadric.cu => geoquadric.inl} | 0 hscgv_uebung_5/lightobject.h | 6 +- .../{lightobject.cpp => lightobject.inl} | 18 +- hscgv_uebung_5/ray.cu | 167 ------------------ hscgv_uebung_5/ray.h | 18 +- hscgv_uebung_5/{ray.cpp => ray.inl} | 87 +++++++-- hscgv_uebung_5/raytracer.cpp | 124 ------------- hscgv_uebung_5/raytracer.cu | 114 ++++++++++++ hscgv_uebung_5/raytracer.h | 15 ++ hscgv_uebung_5/src.pro | 12 +- hscgv_uebung_5/vector.h | 94 +++++----- 14 files changed, 319 insertions(+), 408 deletions(-) rename hscgv_uebung_5/{geoobject.cu => geoobject.inl} (62%) rename hscgv_uebung_5/{geoquadric.cu => geoquadric.inl} (100%) rename hscgv_uebung_5/{lightobject.cpp => lightobject.inl} (69%) delete mode 100644 hscgv_uebung_5/ray.cu rename hscgv_uebung_5/{ray.cpp => ray.inl} (59%) delete mode 100644 hscgv_uebung_5/raytracer.cpp diff --git a/hscgv_uebung_5/geoobject.h b/hscgv_uebung_5/geoobject.h index 8242ea8..0e3ecd8 100644 --- a/hscgv_uebung_5/geoobject.h +++ b/hscgv_uebung_5/geoobject.h @@ -33,11 +33,11 @@ class GeoObjectProperties ~GeoObjectProperties() {}; // access methods - Color ambient() const; - Vec3d reflectance() const; - double specular() const; - int specularExp() const; - double mirror() const; + Color __host__ __device__ ambient() const; + Vec3d __host__ __device__ reflectance() const; + double __host__ __device__ specular() const; + int __host__ __device__ specularExp() const; + double __host__ __device__ mirror() const; protected: Color m_ambient; @@ -57,22 +57,23 @@ class GeoObject // CONSTRUCTORS GeoObject() {}; virtual ~GeoObject() {}; - virtual Vec3d getNormal(const Vec3d &v) const = 0; - virtual double intersect(const Ray &r) const = 0; + virtual Vec3d __host__ __device__ getNormal(const Vec3d &v) const = 0; + virtual double __host__ __device__ intersect(const Ray &r) const = 0; // access methods - virtual Color ambient() const; - virtual Vec3d reflectance() const; - virtual double specular() const; - virtual double specularExp() const; + virtual Color __host__ __device__ ambient() const; + virtual Vec3d __host__ __device__ reflectance() const; + virtual double __host__ __device__ specular() const; + virtual double __host__ __device__ specularExp() const; - virtual double mirror() const; + virtual double __host__ __device__ mirror() const; // config method - virtual void setProperties(GeoObjectProperties *p); + virtual void __host__ __device__ setProperties(GeoObjectProperties *p); protected: GeoObjectProperties *m_properties; }; +#include "geoobject.inl" #endif /* GEOOBJECT_HH */ diff --git a/hscgv_uebung_5/geoobject.cu b/hscgv_uebung_5/geoobject.inl similarity index 62% rename from hscgv_uebung_5/geoobject.cu rename to hscgv_uebung_5/geoobject.inl index dbfe970..fd5ba68 100644 --- a/hscgv_uebung_5/geoobject.cu +++ b/hscgv_uebung_5/geoobject.inl @@ -8,31 +8,31 @@ /* member access functions for GeoObjectProperties */ -Color __host__ __device__ +inline Color __host__ __device__ GeoObjectProperties::ambient() const { return m_ambient; } -Vec3d __host__ __device__ +inline Vec3d __host__ __device__ GeoObjectProperties::reflectance() const { return m_reflectance; } -double __host__ __device__ +inline double __host__ __device__ GeoObjectProperties::specular() const { return m_specular; } -int __host__ __device__ +inline int __host__ __device__ GeoObjectProperties::specularExp() const { return m_specularExp; } -double __host__ __device__ +inline double __host__ __device__ GeoObjectProperties::mirror() const { return m_mirror; @@ -40,57 +40,57 @@ GeoObjectProperties::mirror() const /* access functions for the properties of a GeoObject */ -Color __host__ __device__ +inline Color __host__ GeoObject::ambient() const { if (m_properties) return m_properties->ambient(); else - std::cerr<<"WARNING: properties not set"<reflectance(); else - std::cerr<<"WARNING: properties not set"<specular(); else - std::cerr<<"WARNING: properties not set"<specularExp(); else - std::cerr<<"WARNING: properties not set"<mirror(); else - std::cerr<<"WARNING: properties not set"<direction = dir.getNormalized(); @@ -24,27 +24,27 @@ LightObject::LightObject(const Vec3d& dir, const Color& col) // Description: // Destructor. -LightObject::~LightObject() +inline LightObject::~LightObject() { delete m_properties; } -Color +inline __device__ __host__ Color LightObject::color() const { if (m_properties) return Color(m_properties->color); - else - std::cerr<<"WARNING: properties not set"<direction); - else - std::cerr<<"WARNING: properties not set"< - * Changed by Martin Aumueller - */ - -#include -#include -#include -#include "ray.h" -#include "param.h" -#include "types.h" - -// Description: -// Constructor. -Ray::Ray( ) -: m_origin() -, m_direction() -, m_depth(0) -, m_objList(NULL) -, m_lightList(NULL) -{ - ERR("don't use uninitialized rays !"); -} - -// Description: -// Copy-Constructor. -Ray::Ray( const Ray &r ) -{ - // copy it ! initialization is not enough ! - m_origin = r.m_origin; - m_direction = r.m_direction; - m_depth = r.m_depth; - m_objList = r.m_objList; - m_lightList = r.m_lightList; -} - -// Description: -// Constructor with explicit parameters -Ray::Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll) -{ - // copy it ! initialization is not enough ! - m_origin = o; - m_direction = d; - m_depth = i; - m_objList = &ol; - m_lightList = ≪ -} - -// Description: -// Destructor. -Ray::~Ray( ) -{ -} - -// Description: -// Determine color of this ray by tracing through the scene -const Color __device__ -Ray::shade() const -{ - for (int i=0; i<5; i++) { - GeoObject *closest = NULL; - double tMin = DBL_MAX; - - // find closest object that intersects - for (std::vector::iterator iter = m_objList->begin(); - iter != m_objList->end(); - iter++) { - - double t = (*iter)->intersect(*this); - if (0.0 < t && t < tMin) { - tMin = t; - closest = (*iter); - } - } - - // no object hit -> ray goes to infinity - if (closest == NULL) { - if (m_depth == 0) { - return g_scene.picture.background; // background color - } - else { - return Color(0.0); // black - } - } - else { - // reflection - Vec3d intersectionPosition(m_origin + (m_direction * tMin)); - Vec3d normal(closest->getNormal(intersectionPosition)); - Ray reflectedRay(intersectionPosition, - m_direction.getReflectedAt(normal).getNormalized(), - m_depth+1,*m_objList,*m_lightList); - Color currentColor(0.0); - - // calculate lighting - for (std::vector::iterator iter = m_lightList->begin(); - iter != m_lightList->end(); - iter++) { - - // where is the lightsource ? - Ray rayoflight(intersectionPosition, (*iter)->direction(), 0, *m_objList, *m_lightList); - bool something_intersected = false; - - // where are the objects ? - for (std::vector::iterator iter2 = m_objList->begin(); - iter2 != m_objList->end(); - iter2++) { - - double t = (*iter2)->intersect(rayoflight); - if (t > 0.0) { - something_intersected = true; - break; - } - - } // for all obj - - // is it visible ? - if (! something_intersected) - currentColor += shadedColor((*iter), reflectedRay, normal, closest); - - } // for all lights - - // could be right... - currentColor *= closest->mirror(); - } - } - return Color(currentColor); -} - -// Description: -// Determine color contribution of a lightsource -const Color __device__ __host__ -Ray::shadedColor(LightObject *light, const Ray &reflectedRay, const Vec3d &normal, GeoObject *obj) const -{ - double ldot = light->direction() | normal; - Color reflectedColor = Color(0.0); - - // lambertian reflection model - if (ldot > 0.0) - reflectedColor += obj->reflectance() * (light->color() * ldot); - - // updated with ambient lightning as in: - // [GENERALISED AMBIENT REFLECTION MODELS FOR LAMBERTIAN AND PHONG SURFACES, Xiaozheng Zhang and Yongsheng Gao] - reflectedColor += obj->ambient() * g_scene.ambience; - - // specular part - double spec = reflectedRay.direction() | light->direction(); - if (spec > 0.0) { - spec = obj->specular() * pow(spec, obj->specularExp()); - reflectedColor += light->color() * spec; - } - - return Color(reflectedColor); -} - -Vec3d __device__ __host__ -Ray::origin() const -{ - return Vec3d(m_origin); -} - -Vec3d __device__ __host__ -Ray::direction() const -{ - return Vec3d(m_direction); -} - diff --git a/hscgv_uebung_5/ray.h b/hscgv_uebung_5/ray.h index 1151d3d..e49e304 100644 --- a/hscgv_uebung_5/ray.h +++ b/hscgv_uebung_5/ray.h @@ -23,19 +23,20 @@ class Ray { public: // CONSTRUCTORS - Ray(); - Ray(const Ray &r); - Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll); - ~Ray(); + __host__ __device__ Ray(); + __host__ __device__ Ray(const Ray &r); + __host__ __device__ Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll); + __host__ __device__ ~Ray(); - const Color shade() const; + const Color __host__ shade() const; + const Color __device__ cudaShade(GeoObject* m_objListCuda, int objListSize, LightObject *m_lightListCuda, int lightListSize) const; // access methods - Vec3d origin() const; - Vec3d direction() const; + Vec3d __host__ __device__ origin() const; + Vec3d __host__ __device__ direction() const; protected: - const Color shadedColor(LightObject *light, const Ray &reflectedray, const Vec3d &normal, GeoObject *obj) const; + const Color __host__ __device__ shadedColor(LightObject *light, const Ray &reflectedray, const Vec3d &normal, GeoObject *obj) const; private: Vec3d m_origin; @@ -45,5 +46,6 @@ class Ray std::vector *m_objList; std::vector *m_lightList; }; +#include "ray.inl" #endif /* RAY_HH */ diff --git a/hscgv_uebung_5/ray.cpp b/hscgv_uebung_5/ray.inl similarity index 59% rename from hscgv_uebung_5/ray.cpp rename to hscgv_uebung_5/ray.inl index 99539e8..c3d47ed 100644 --- a/hscgv_uebung_5/ray.cpp +++ b/hscgv_uebung_5/ray.inl @@ -13,19 +13,18 @@ // Description: // Constructor. -Ray::Ray( ) +inline Ray::Ray( ) : m_origin() , m_direction() , m_depth(0) , m_objList(NULL) , m_lightList(NULL) { - ERR("don't use uninitialized rays !"); } // Description: // Copy-Constructor. -Ray::Ray( const Ray &r ) +inline Ray::Ray( const Ray &r ) { // copy it ! initialization is not enough ! m_origin = r.m_origin; @@ -37,7 +36,7 @@ Ray::Ray( const Ray &r ) // Description: // Constructor with explicit parameters -Ray::Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll) +inline Ray::Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll) { // copy it ! initialization is not enough ! m_origin = o; @@ -49,13 +48,13 @@ Ray::Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector // Description: // Destructor. -Ray::~Ray( ) +inline Ray::~Ray( ) { } // Description: // Determine color of this ray by tracing through the scene -const Color +inline const Color __host__ Ray::shade() const { GeoObject *closest = NULL; @@ -126,10 +125,77 @@ Ray::shade() const return Color(currentColor); } } +// Description: +// Determine color of this ray by tracing through the scene +inline const Color __device__ +Ray::cudaShade(GeoObject* m_objListCuda, int objListSize, LightObject* m_lightListCuda, int lightListSize) const +{ + Color currentColor(0.0); + for (int i=0; i<5; i++) { + GeoObject *closest = NULL; + double tMin = DBL_MAX; + + // find closest object that intersects + for (int j=0; j ray goes to infinity + if (closest == NULL) { + if (m_depth == 0) { +// return g_sceneCuda.picture.background; // background color + } + else { + return Color(0.0); // black + } + } + else { + // reflection + Vec3d intersectionPosition(m_origin + (m_direction * tMin)); + Vec3d normal(closest->getNormal(intersectionPosition)); + Ray reflectedRay(intersectionPosition, + m_direction.getReflectedAt(normal).getNormalized(), + m_depth+1,*m_objList,*m_lightList); + + // calculate lighting + for (int j=0; j 0.0) { + something_intersected = true; + break; + } + + } // for all obj + + // is it visible ? + if (! something_intersected) + currentColor += shadedColor(&m_lightListCuda[j], reflectedRay, normal, closest); + + } // for all lights + + // could be right... + currentColor *= closest->mirror(); + } + } + return Color(currentColor); +} // Description: // Determine color contribution of a lightsource -const Color +inline const Color __device__ __host__ Ray::shadedColor(LightObject *light, const Ray &reflectedRay, const Vec3d &normal, GeoObject *obj) const { double ldot = light->direction() | normal; @@ -141,7 +207,7 @@ Ray::shadedColor(LightObject *light, const Ray &reflectedRay, const Vec3d &norma // updated with ambient lightning as in: // [GENERALISED AMBIENT REFLECTION MODELS FOR LAMBERTIAN AND PHONG SURFACES, Xiaozheng Zhang and Yongsheng Gao] - reflectedColor += obj->ambient() * g_scene.ambience; +// reflectedColor += obj->ambient() * g_sceneCuda.ambience; // specular part double spec = reflectedRay.direction() | light->direction(); @@ -153,14 +219,15 @@ Ray::shadedColor(LightObject *light, const Ray &reflectedRay, const Vec3d &norma return Color(reflectedColor); } -Vec3d +inline Vec3d __device__ __host__ Ray::origin() const { return Vec3d(m_origin); } -Vec3d +inline Vec3d __device__ __host__ Ray::direction() const { return Vec3d(m_direction); } + diff --git a/hscgv_uebung_5/raytracer.cpp b/hscgv_uebung_5/raytracer.cpp deleted file mode 100644 index 5bf2907..0000000 --- a/hscgv_uebung_5/raytracer.cpp +++ /dev/null @@ -1,124 +0,0 @@ -/* ******** Programmierpraktikum Computergrafik (CGP) ************** - * Aufgabe 1 - "Lichtblick" - * Created by Peter Kipfer - * Changed by Martin Aumueller - */ - -#include "raytracer.h" - -// Description: -// the connection to the parser -extern int openSceneFile(const char *s); -extern int closeSceneFile(void); -extern int inputparse(void); -extern int cleanUp(void); - -//! get scene description from the parser -/*! - Get the scene description from the parser. - \param inputfile Filename of the scene - */ -void -ReadScene(const char *inputfile) -{ - openSceneFile(inputfile); - inputparse(); - closeSceneFile(); -} - -// scene storage -//! the list of geometric objects in the scene -extern std::vector g_objectList; -//! the list of light sources in the scene -extern std::vector g_lightList; -//! the list of surface properties of the objects -extern std::vector g_propList; - -Raytracer::Raytracer(): - m_isFileLoaded(false), - m_filename(0), - m_antialiasing(false) {} -Raytracer::Raytracer(bool antialiasing): - m_isFileLoaded(false), - m_filename(0), - m_antialiasing(antialiasing) {} -Raytracer::Raytracer(const char *filename, bool antialiasing): - m_filename(filename), - m_antialiasing(antialiasing) -{ - // parse the input file - ReadScene(m_filename); - m_isFileLoaded = true; -} -Raytracer::~Raytracer() { - // clean up - cleanUp(); -} - -void -Raytracer::start(float *renderedScene, int xRes, int yRes) { - // setup viewport, its origin is bottom left - // setup camera coordsys - Vec3d eye_dir = (g_scene.view.lookat - g_scene.view.eyepoint).getNormalized(); - Vec3d eye_right = (eye_dir^g_scene.view.up).getNormalized(); - Vec3d eye_up = eye_dir^eye_right*-1; - - // calculatehe dimensions of the viewport using the scene's camera - float height = 2 * tan(M_PI/180 * .5 * g_scene.view.fovy); - float width = height * g_scene.view.aspect; - - // compute delta steps in each direction - Vec3d deltaX = eye_right * (width / xRes); - Vec3d deltaY = eye_up * (height / yRes); - - // this should be bottom left - Vec3d bottomLeft = g_scene.view.eyepoint + eye_dir - deltaX*xRes/2 - deltaY*yRes/2; - - // normal ray tracing: the color of the center of a pixel is computed - #pragma omp parallel for schedule(dynamic) collapse(2) - for (int sy=yRes ; sy > 0 ; --sy) { - for (int sx=0 ; sx < xRes ; ++sx) { - // the center of the pixel we are looking at right now - Vec3d point = bottomLeft + deltaX*sx + deltaY*sy + deltaX/2 + deltaY/2; - - // the direction of our look - Vec3d dir = point - g_scene.view.eyepoint; - - // create ray from view.eyepoint to view.lookat - Ray theRay(g_scene.view.eyepoint,dir.getNormalized(),0,g_objectList,g_lightList); - - // compute the color - Color col = theRay.shade(); - - // in case we are using antialiasing, calculate the color of this pixel by averaging - if (m_antialiasing) { - // scale the midpoint color since we are going to use 5 points to average our color - col *= 0.2; - - // besides taking shooting a ray through the midpoint 'o', we calculate - // the pixels color by shooting 4 more rays through the points 'x' - // and averaging their values - // ------- - // | x x | - // | o | - // | x x | - // -------- - for(float dx = -1/4.; dx <= 1/4.; dx+=1/2.) { - for(float dy = -1/4.; dy <= 1/4.; dy+=1/2.) { - Vec3d superSamplePoint = point + deltaX*dx + deltaY*dy; - Vec3d superSampleDir = superSamplePoint - g_scene.view.eyepoint; - - // create ray from view.eyepoint to view.lookat - Ray theRay(g_scene.view.eyepoint,superSampleDir.getNormalized(),0,g_objectList,g_lightList); - col += theRay.shade()*0.2;//color, recursive_ray_trace(eye, ray, 0)); - } - } - } - - int index = 3*((sy-1) * xRes + sx); - renderedScene[index + 0] = col[0]; - renderedScene[index + 1] = col[1]; - renderedScene[index + 2] = col[2]; - } // foreach x - } // foreach y -} diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index ad9f0f2..2516c63 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -1,6 +1,48 @@ +/* ******** Programmierpraktikum Computergrafik (CGP) ************** + * Aufgabe 1 - "Lichtblick" + * Created by Peter Kipfer + * Changed by Martin Aumueller + */ + #include "raytracer.h" +#include "omp.h" +//! get scene description from the parser +/*! + Get the scene description from the parser. + \param inputfile Filename of the scene + */ void +ReadScene(const char *inputfile) +{ + openSceneFile(inputfile); + inputparse(); + closeSceneFile(); +} + + +Raytracer::Raytracer(): + m_isFileLoaded(false), + m_filename(0), + m_antialiasing(false) {} +Raytracer::Raytracer(bool antialiasing): + m_isFileLoaded(false), + m_filename(0), + m_antialiasing(antialiasing) {} +Raytracer::Raytracer(const char *filename, bool antialiasing): + m_filename(filename), + m_antialiasing(antialiasing) +{ + // parse the input file + ReadScene(m_filename); + m_isFileLoaded = true; +} +Raytracer::~Raytracer() { + // clean up + cleanUp(); +} + +__host__ void Raytracer::start(float *renderedScene, int xRes, int yRes) { // setup viewport, its origin is bottom left // setup camera coordsys @@ -20,6 +62,78 @@ Raytracer::start(float *renderedScene, int xRes, int yRes) { Vec3d bottomLeft = g_scene.view.eyepoint + eye_dir - deltaX*xRes/2 - deltaY*yRes/2; // normal ray tracing: the color of the center of a pixel is computed +#ifndef __CUDA_ARCH__ + fprintf(stderr, "%d threads on board \n", omp_get_num_threads()); + #pragma omp parallel for schedule(dynamic) collapse(2) +#endif + for (int sy=yRes ; sy > 0 ; --sy) { + for (int sx=0 ; sx < xRes ; ++sx) { + // the center of the pixel we are looking at right now + Vec3d point = bottomLeft + deltaX*sx + deltaY*sy + deltaX/2 + deltaY/2; + + // the direction of our look + Vec3d dir = point - g_scene.view.eyepoint; + + // create ray from view.eyepoint to view.lookat + Ray theRay(g_scene.view.eyepoint,dir.getNormalized(),0,g_objectList,g_lightList); + + // compute the color + Color col = theRay.shade(); + + // in case we are using antialiasing, calculate the color of this pixel by averaging + if (m_antialiasing) { + // scale the midpoint color since we are going to use 5 points to average our color + col *= 0.2; + + // besides taking shooting a ray through the midpoint 'o', we calculate + // the pixels color by shooting 4 more rays through the points 'x' + // and averaging their values + // ------- + // | x x | + // | o | + // | x x | + // -------- + for(float dx = -1/4.; dx <= 1/4.; dx+=1/2.) { + for(float dy = -1/4.; dy <= 1/4.; dy+=1/2.) { + Vec3d superSamplePoint = point + deltaX*dx + deltaY*dy; + Vec3d superSampleDir = superSamplePoint - g_scene.view.eyepoint; + + // create ray from view.eyepoint to view.lookat + Ray theRay(g_scene.view.eyepoint,superSampleDir.getNormalized(),0,g_objectList,g_lightList); + col += theRay.shade()*0.2;//color, recursive_ray_trace(eye, ray, 0)); + } + } + } + + int index = 3*((sy-1) * xRes + sx); + renderedScene[index + 0] = col[0]; + renderedScene[index + 1] = col[1]; + renderedScene[index + 2] = col[2]; + } // foreach x + } // foreach y +} + + +void +Raytracer::startCuda(float *renderedScene, int xRes, int yRes) { + // setup viewport, its origin is bottom left + // setup camera coordsys + Vec3d eye_dir = (g_scene.view.lookat - g_scene.view.eyepoint).getNormalized(); + Vec3d eye_right = (eye_dir^g_scene.view.up).getNormalized(); + Vec3d eye_up = eye_dir^eye_right*-1; + + // calculatehe dimensions of the viewport using the scene's camera + float height = 2 * tan(M_PI/180 * .5 * g_scene.view.fovy); + float width = height * g_scene.view.aspect; + + // compute delta steps in each direction + Vec3d deltaX = eye_right * (width / xRes); + Vec3d deltaY = eye_up * (height / yRes); + + // this should be bottom left + Vec3d bottomLeft = g_scene.view.eyepoint + eye_dir - deltaX*xRes/2 - deltaY*yRes/2; + + // normal ray tracing: the color of the center of a pixel is computed for (int sy=yRes ; sy > 0 ; --sy) { for (int sx=0 ; sx < xRes ; ++sx) { // the center of the pixel we are looking at right now diff --git a/hscgv_uebung_5/raytracer.h b/hscgv_uebung_5/raytracer.h index d1859b0..35f483a 100644 --- a/hscgv_uebung_5/raytracer.h +++ b/hscgv_uebung_5/raytracer.h @@ -28,10 +28,25 @@ class Raytracer Raytracer(const char* filename, bool antialiasing); ~Raytracer(); void start(float* renderedScene, int xRes, int yRes); + void startCuda(float* renderedScene, int xRes, int yRes); bool m_isFileLoaded; private: const char *m_filename; bool m_antialiasing; }; +// Description: +// the connection to the parser +extern int openSceneFile(const char *s); +extern int closeSceneFile(void); +extern int inputparse(void); +extern int cleanUp(void); +// scene storage +//! the list of geometric objects in the scene +extern std::vector g_objectList; +//! the list of light sources in the scene +extern std::vector g_lightList; +//! the list of surface properties of the objects +extern std::vector g_propList; + #endif // RAYTRACER_H diff --git a/hscgv_uebung_5/src.pro b/hscgv_uebung_5/src.pro index 9cdeebd..eacb14b 100644 --- a/hscgv_uebung_5/src.pro +++ b/hscgv_uebung_5/src.pro @@ -6,18 +6,18 @@ HEADERS= geoobject.h geoquadric.h lightobject.h ray.h types.h vector.h pa ApplicationWindow.h \ GLFrame.h \ raytracer.h -SOURCES= lightobject.cpp main.cpp ray.cpp \ geopolygon.cpp \ +SOURCES= main.cpp \ ApplicationWindow.cpp \ GLFrame.cpp \ - raytracer.cpp + ray.inl \ + geoobject.inl \ + geoquadric.inl \ + lightobject.inl win32:SOURCES*= xgetopt.cpp YACCSOURCES= input.y CUDA_SOURCES += \ - raytracer.cu \ - ray.cu \ - geoobject.cu \ - geoquadric.cu + raytracer.cu QMAKE_YACC = bison QMAKE_YACCFLAGS = -y -d diff --git a/hscgv_uebung_5/vector.h b/hscgv_uebung_5/vector.h index 81fa741..4532572 100644 --- a/hscgv_uebung_5/vector.h +++ b/hscgv_uebung_5/vector.h @@ -43,129 +43,129 @@ class Vector { public: //! Constructor - Vector(void ); + __host__ __device__ Vector(void ); //! Copy-Constructor - Vector(const Vector &v ); + __host__ __device__ Vector(const Vector &v ); //! construct a Vector from one Scalar - Vector(Scalar); + __host__ __device__ Vector(Scalar); //! construct a Vector from three Scalars - Vector(Scalar, Scalar, Scalar); + __host__ __device__ Vector(Scalar, Scalar, Scalar); //! Assignment operator - const Vector& operator= (const Vector& v); + __host__ __device__ const Vector& operator= (const Vector& v); //! Assignment operator - const Vector& operator= (Scalar s); + __device__ __host__ const Vector& operator= (Scalar s); //! Assign and add operator - const Vector& operator+= (const Vector& v); + __device__ __host__ const Vector& operator+= (const Vector& v); //! Assign and add operator - const Vector& operator+= (Scalar s); + __device__ __host__ const Vector& operator+= (Scalar s); //! Assign and sub operator - const Vector& operator-= (const Vector& v); + __device__ __host__ const Vector& operator-= (const Vector& v); //! Assign and sub operator - const Vector& operator-= (Scalar s); + __device__ __host__ const Vector& operator-= (Scalar s); //! Assign and mult operator - const Vector& operator*= (const Vector& v); + __device__ __host__ const Vector& operator*= (const Vector& v); //! Assign and mult operator - const Vector& operator*= (Scalar s); + __device__ __host__ const Vector& operator*= (Scalar s); //! Assign and div operator - const Vector& operator/= (const Vector& v); + __device__ __host__ const Vector& operator/= (const Vector& v); //! Assign and div operator - const Vector& operator/= (Scalar s); + __device__ __host__ const Vector& operator/= (Scalar s); //! output operator - friend std::ostream& operator<< <>( std::ostream& os, const Vector& inst ); + __device__ __host__ friend std::ostream& operator<< <>( std::ostream& os, const Vector& inst ); //! input operator friend std::istream& operator>> <>( std::istream& is, Vector& inst ); //! unary operator - Vector operator- () const; + __host__ __device__ Vector operator- () const; //! binary operator add - Vector operator+ (const Vector&) const; + __host__ __device__ Vector operator+ (const Vector&) const; //! binary operator add - Vector operator+ (Scalar) const; + __host__ __device__ Vector operator+ (Scalar) const; //! binary operator sub - Vector operator- (const Vector&) const; + __host__ __device__ Vector operator- (const Vector&) const; //! binary operator sub - Vector operator- (Scalar) const; + __host__ __device__ Vector operator- (Scalar) const; //! binary operator mult - Vector operator* (const Vector&) const; + __host__ __device__ Vector operator* (const Vector&) const; //! binary operator mult - Vector operator* (Scalar) const; + __host__ __device__ Vector operator* (Scalar) const; //! binary operator div - Vector operator/ (const Vector&) const; + __host__ __device__ Vector operator/ (const Vector&) const; //! binary operator div - Vector operator/ (Scalar) const; + __host__ __device__ Vector operator/ (Scalar) const; //! equality test operator - bool operator== (const Vector&) const; + __host__ __device__ bool operator== (const Vector&) const; //! inequality test operator - bool operator!= (const Vector&) const; + __host__ __device__ bool operator!= (const Vector&) const; //! scalar product - Scalar operator|( const Vector &v ) const; + __host__ __device__ Scalar operator|( const Vector &v ) const; //! norm - Scalar getNorm() const; + __host__ __device__ Scalar getNorm() const; //! normalize - Vector getNormalized() const; + __host__ __device__ Vector getNormalized() const; //! normalize and norm - Scalar normalize(); + __host__ __device__ Scalar normalize(); //! Vector product - Vector operator^( const Vector &v ) const; + __host__ __device__ Vector operator^( const Vector &v ) const; //! Projection normal to a vector - Vector getOrthogonalVector() const; + __host__ __device__ Vector getOrthogonalVector() const; //! Project into a plane - const Vector& projectNormalTo(const Vector &v); + __host__ __device__ const Vector& projectNormalTo(const Vector &v); //! Reflect at normal - Vector getReflectedAt(const Vector& n) const; + __host__ __device__ Vector getReflectedAt(const Vector& n) const; //! Refract at normal - Vector getRefractedAt(const Vector& n, + __host__ __device__ Vector getRefractedAt(const Vector& n, double index, bool& totalReflection) const; //! minimize - const Vector &minimize(const Vector &); + __host__ __device__ const Vector &minimize(const Vector &); //! maximize - const Vector &maximize(const Vector &); + __host__ __device__ const Vector &maximize(const Vector &); //! access operator - Scalar& operator[](unsigned int i); + __host__ __device__ Scalar& operator[](unsigned int i); //! access operator - const Scalar& operator[](unsigned int i) const; + __host__ __device__ const Scalar& operator[](unsigned int i) const; //! Get minimal vector length - static Scalar getEpsilon( ); + __host__ __device__ static Scalar getEpsilon( ); protected: @@ -252,7 +252,7 @@ typedef Vector Color; /*! Constructor. */ -template +template __device__ __host__ Vector::Vector( void ) { #ifdef VECTOR_INIT_ZERO @@ -265,7 +265,7 @@ Vector::Vector( void ) /*! Copy-Constructor. */ -template +template __device__ __host__ Vector::Vector( const Vector &v ) { m_value[0] = v.m_value[0]; @@ -281,7 +281,7 @@ Vector::Vector( const Vector &v ) \param s The m_value to set \return The new Vector */ -template +template __device__ __host__ Vector::Vector(Scalar s ) { m_value[0]= s; @@ -297,7 +297,7 @@ Vector::Vector(Scalar s ) \param s3 The value for the third Vector component \return The new Vector */ -template +template __device__ __host__ Vector::Vector(Scalar s1, Scalar s2, Scalar s3) { m_value[0]= s1; @@ -311,7 +311,7 @@ Vector::Vector(Scalar s1, Scalar s2, Scalar s3) \param v Second vector to compute the product with \return A new Vector with the product values */ -template +template __device__ __host__ Vector Vector::operator^( const Vector &v ) const { @@ -326,7 +326,7 @@ Vector::operator^( const Vector &v ) const \param v Vector with values to be copied \return Reference to self */ -template +template __device__ __host__ const Vector& Vector::operator=( const Vector &v ) { @@ -342,7 +342,7 @@ Vector::operator=( const Vector &v ) \param s The value to copy \return Reference to self */ -template +template __device__ __host__ const Vector& Vector::operator=(Scalar s) { From 8bb3e6e55a4ea52d2d85ac4aa96d01e103e503fc Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 18 May 2015 18:05:58 +0200 Subject: [PATCH 19/87] typo --- hscgv_uebung_5/vector.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hscgv_uebung_5/vector.h b/hscgv_uebung_5/vector.h index 4532572..4e0be9e 100644 --- a/hscgv_uebung_5/vector.h +++ b/hscgv_uebung_5/vector.h @@ -13,7 +13,7 @@ #include #include -#include +#include "types.h" //! should we initialize with zero ? /*! From bf5478db71930e2f50272410bb0e62a362f5fa17 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 18 May 2015 18:06:22 +0200 Subject: [PATCH 20/87] remove polygon functionality --- hscgv_uebung_5/geopolygon.cpp | 202 ---------------------------------- hscgv_uebung_5/geopolygon.h | 58 ---------- hscgv_uebung_5/input.y | 132 +--------------------- hscgv_uebung_5/src.pro | 2 - 4 files changed, 1 insertion(+), 393 deletions(-) delete mode 100644 hscgv_uebung_5/geopolygon.cpp delete mode 100644 hscgv_uebung_5/geopolygon.h diff --git a/hscgv_uebung_5/geopolygon.cpp b/hscgv_uebung_5/geopolygon.cpp deleted file mode 100644 index 234cb7b..0000000 --- a/hscgv_uebung_5/geopolygon.cpp +++ /dev/null @@ -1,202 +0,0 @@ -/* - * Created by Jeronim Morina - * - * File: geopolygon.cpp - * Declaration of polygonal surfaces - */ -#include -#include -#include "geopolygon.h" - -// Description: -// Constructor. -GeoPolygon::GeoPolygon() -{ - m_properties = NULL; -} - -// Description: -// add a new polygonal surface with a list of all vertices -GeoPolygon::GeoPolygon(std::vector vertices, std::vector > polygons) -: m_vertices(vertices), m_polygons(polygons) -{ - m_properties = NULL; -} - -// Description: -// Destructor. -GeoPolygon::~GeoPolygon() -{ - m_polygons.clear(); - m_vertices.clear(); -} - -// Description: -// Return the normal vector at point v of this surface -Vec3d -GeoPolygon::getNormal(const Vec3d &v) const -{ - // find out where v lies - for (unsigned poly=0; polytesty) != (m_vertices[indices[j]][cVec[1]]>testy)) && - (testx < (m_vertices[indices[j]][cVec[0]]-m_vertices[indices[i]][cVec[0]])* (testy-m_vertices[indices[i]][cVec[1]]) / (m_vertices[indices[j]][cVec[1]]-m_vertices[indices[i]][cVec[1]]) + m_vertices[indices[i]][cVec[0]]) ) - c = !c; - } - return c; -} - - -// Description: -// Compute intersection point on this surface -// return distance to nearest intersection found or -1 -double -GeoPolygon::intersect(const Ray &ray) const -{ - double temp, lastTemp=DBL_MAX; - for (unsigned i=0; i= 0 && temp < lastTemp) lastTemp = temp; - } - return (lastTemp p = m_polygons[nPoly]; - Vec3d edge, vp; - Vec3d origin = ray.origin(); - Vec3d dir = ray.direction(); - double t, si; - - Vec3d norm = getNormal(nPoly); - double nDir = norm | dir; - double d = norm | p[0]; - - - if(nDir != 0.0) - { - // determine where the intersection is - si = norm | origin; - t = (d - si) / nDir; - Vec3d point = origin + dir*t; - if ((t > Vec3d::getEpsilon())) - { - unsigned length = p.size(); - - // see if intersection is on polygon edge - for (unsigned i = 0; i < length - 1; i++) - { - edge = p[i+1] - p[i]; - vp = point - p[i]; - dir = edge^vp; - if((vp | dir) < 0) - { - return -1.; - } - } - - // check the last edge - edge = p[length-1] - p[0]; - vp = point - p[0]; - dir = edge^vp; - if((vp | dir) < 0) - { - return -1.; - } - - // see if intersection lies inside the intersection of all spanned planes - for (unsigned i = 1; i < length - 1; i++) { - edge = p[i+1] - p[i]; - vp = point - p[i]; - dir = p[i-1] - p[i]; - if((vp | dir) < 0) - { - return -1.; - } - - } - - // handle very first plane - edge = p[length-1] - p[0]; - vp = point - p[0]; - dir = p[1] - p[0]; - if((vp | dir) < 0) - { - return -1.; - } - - // and the last plane - edge = p[0] - p[1]; - vp = point - p[1]; - dir = p[2] - p[1]; - if((vp | dir) < 0) - { - return -1.; - } - - // all checks survived, the point v is definitely inside this polygon - return t; - } - } - // nDir is 0! - return -1.; -} diff --git a/hscgv_uebung_5/geopolygon.h b/hscgv_uebung_5/geopolygon.h deleted file mode 100644 index 5d01051..0000000 --- a/hscgv_uebung_5/geopolygon.h +++ /dev/null @@ -1,58 +0,0 @@ -/* - * Created by Jeronim Morina - * - * File: geopolygon.h - * Declaration of polygonal surfaces - */ -#ifndef GEOPOLYGON_H -#define GEOPOLYGON_H - -#include "geoobject.h" -#include "ray.h" - - -//! A polygon surface -/*! This type implements a special type of geometric objects, - polygon surfaces given in a vertex-index-list form - */ -class GeoPolygon : public GeoObject -{ - public: - // CONSTRUCTORS - //! Construct trivial \b polygon. - GeoPolygon(); - //! Construct a polygonal surface consisting of a list of vertices and several connected polygons - GeoPolygon(std::vector vertices, std::vector > polygons); - //! Delete a polygon. - virtual ~GeoPolygon(); - - //! Compute surface normal on the polygon in the point \b v. - /*! since we call this method only upon positive intersection, \b v is indeed a - * point on the surface. This method finds out in which polygon it lies and returns - * its easy to compute normal. In case that \b v lies on an edge or a vertex, - * we take the first found polygon. \b Regard that this should be changed in - * case there are too many artefacts. - */ - virtual Vec3d getNormal(const Vec3d &v) const; - //! Compute intersection of ray with polygon. - virtual double intersect(const Ray &r) const; - - private: - //! All vertices belonging to this polygonal surface - std::vector m_vertices; - //! This polygonal surface consists of several polygons - /*! Each polygon is saved as a list of indices, pointing into \b m_vertices */ - std::vector > m_polygons; - - //! This method returns the normal to nPoly-th polygon - Vec3d getNormal(const int nPoly) const; - //! This method returns the distance to the intersection with the nPoly-th polygon - /*! In case that no intersection is found with this polygon we return -1 - * otherwise the distance t */ - double intersect(const Ray &ray, const int nPoly) const; - //! PNPOLY - Point Inclusion in Polygon Test - /*! Copyright (c) 1970-2003, Wm. Randolph Franklin */ - int pnpoly(int nvert, int *cVec, const int *indices, float testx, float testy) const; -}; - -#endif // GEOPOLYGON_H diff --git a/hscgv_uebung_5/input.y b/hscgv_uebung_5/input.y index 0dd0cda..3f08614 100644 --- a/hscgv_uebung_5/input.y +++ b/hscgv_uebung_5/input.y @@ -23,7 +23,6 @@ #include "vector.h" #include "param.h" #include "geoquadric.h" -#include "geopolygon.h" #include "lightobject.h" #include "input_yacc.h" @@ -44,7 +43,6 @@ static double currentSpecularExp = 0; static double currentMirror = 0; static std::vector vertexList; static std::vector indexList; -static std::vector > polygonList; /* * table of the keywords that we use above @@ -61,11 +59,8 @@ keywords[] = { { "aspect", ASPECT }, { "numsurfaces", NUMSURFACES }, { "object", OBJECT }, - { "poly", POLY }, - { "polygon", POLY }, /* equivalent spelling */ { "numvertices", NUMVERTICES }, { "vertex", VERTEX }, - { "numpolygons", NUMPOLYGONS }, { "quadric", QUADRIC }, { "numproperties", NUMPROPERTIES }, { "ambient", AMBIENT }, @@ -347,7 +342,6 @@ static int propertyCounter = 0; static int objectCounter = 0; static int lightCounter = 0; static int vertexCounter = 0; -static int polyCounter = 0; static int indexCounter = 0; /* @@ -358,7 +352,6 @@ static int nsurfaces = 0; static int nobjs = 0; static int nlights = 0; static int nvertices = 0; -static int npolys = 0; static int nindices = 0; /* @@ -601,10 +594,9 @@ surfaces | one_surface ; -/* we know quadric surfaces and polygonal surfaces */ +/* we know quadric surfaces*/ one_surface : quadric_surface - | polygon_surface ; /* a quadric surface is given by the 10 parameters for @@ -624,128 +616,6 @@ quadric_surface } ; -/* The first step for creating a polygon surface is building up - * a list of vertices. Then a list of polygons is created via - * indices into this vertex list. */ -polygon_surface - : OBJECT POLY - { - if ( surfaceCounter >= nsurfaces ) { - yywarn("too many surfaces specified (ignoring)"); - } else { - TOUT((stderr,"object poly\n")); - } - } -vertex_section polygon_section - { - if ( surfaceCounter < nsurfaces ) { - surfaceCounter++; - /* we completed parsing an entire polygonal surface, - now we can create a geoObject out of the input */ - g_objectList.push_back(new GeoPolygon(vertexList, polygonList)); - TOUT((stderr,"We just added a new polygonal surface\n")); - } - } -; - -/* this starts a new list of vertices: clear the old one */ -vertex_section - : NUMVERTICES index - { TOUT((stderr,"numvertices %d\n", $2)); - nvertices = $2; - TOUT((stderr,"Begin parsing the vertices. Clear old vertexList.\n")); - vertexList.clear(); - } - -/* completed parsing vertices, do nothing! */ -vertices - { TOUT((stderr,"Completed parsing all vertices.\n")); } -; - -/* vertices follow each other */ -vertices - : vertices one_vertex - | one_vertex -; - -/* parse one vertex and store it into our temporary vertex list */ -one_vertex - : VERTEX realVal realVal realVal - { TOUT((stderr,"vertex %f %f %f\n", $2, $3, $4)); - if(vertexCounter >= nvertices) { - yywarn("too many vertices specified (ignoring)"); - } else { - vertexList.push_back(Vec3d($2,$3,$4)); - } - } -; - -/* this starts a new polygon surface: clear the old list */ -polygon_section - : NUMPOLYGONS index - { TOUT((stderr,"numpolygons %d\n", $2)); - npolys = $2; - polygonList.clear(); - polyCounter = 0; - } - -/* completed parsing indices, do nothing! */ -polygons - { TOUT((stderr,"Completed parsing all indices\n")); - } -; - -/* a polygon object consists of several polygons */ -polygons - : polygons one_polygon - | one_polygon -; - -/* clear the old polygon as a new one is started parsing */ -one_polygon - : POLY index - { TOUT((stderr,"polygon %d ", $2)); - nindices = $2; - indexList.clear(); - indexCounter = 0; - } - -/* add polygon to the list of polygons belonging to this surface */ -indices - { TOUT((stderr,"Add the whole indexList to our polygons\n")); - polygonList.push_back(indexList); - polyCounter++; - } -; - -/* read the indices into the indexList that describes one polygon */ -indices - : indices index - { TOUT((stderr," %d", $2)); - if($2<=0 || $2>nvertices) { - yywarn("vertex index out of range (ignoring)"); - } else if(indexCounter > nindices) { - yywarn("too many vertices specified (ignoring)"); - } else { - /* add the first index ($2) of a new polygon in the polygon list (polygons.push_back()) - of the current polygon surface structure (g_objectList.back()) */ - indexList.push_back($2); - indexCounter++; - } - } - | index - { TOUT((stderr," %d", $1)); - if($1<=0 || $1>nvertices) { - yywarn("vertex index out of range (ignoring)"); - } else if(indexCounter > nindices) { - yywarn("too many vertices specified (ignoring)"); - } else { - /* add another index to the current polygon (there should be at least 3) */ - indexList.push_back($1); - indexCounter++; - } - } -; /* parse object properties */ property_section diff --git a/hscgv_uebung_5/src.pro b/hscgv_uebung_5/src.pro index eacb14b..08307a7 100644 --- a/hscgv_uebung_5/src.pro +++ b/hscgv_uebung_5/src.pro @@ -2,11 +2,9 @@ TARGET= raytrace QT *= opengl HEADERS= geoobject.h geoquadric.h lightobject.h ray.h types.h vector.h param.h \ - geopolygon.h \ ApplicationWindow.h \ GLFrame.h \ raytracer.h - geopolygon.cpp \ SOURCES= main.cpp \ ApplicationWindow.cpp \ GLFrame.cpp \ From f5369553096bb7cbf7ed31d09bd1e9281188f57c Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 18 May 2015 18:06:27 +0200 Subject: [PATCH 21/87] typo --- hscgv_uebung_5/GLFrame.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index b1d53fe..7a484d9 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -13,7 +13,7 @@ #ifdef __APPLE__ #include #include -#elif +#else #include #include #endif From 08d6308af4016167f5b6f33be80c78563bc2556e Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 18 May 2015 18:06:40 +0200 Subject: [PATCH 22/87] setup qmake lflags for openmp --- hscgv_uebung_5/src.pro | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hscgv_uebung_5/src.pro b/hscgv_uebung_5/src.pro index 08307a7..695fdd0 100644 --- a/hscgv_uebung_5/src.pro +++ b/hscgv_uebung_5/src.pro @@ -26,7 +26,7 @@ CONFIG += debug CONFIG -= release QMAKE_CXXFLAGS += -W -Wall -fopenmp -#QMAKE_LFLAGS += -fopenmp +QMAKE_LFLAGS += -fopenmp FORMS += \ ApplicationWindow.ui From 7dd1d565b276f73f346c5cf5293e41a5a8443aec Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 22 May 2015 17:58:39 +0200 Subject: [PATCH 23/87] update application name --- hscgv_uebung_5/ApplicationWindow.ui | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hscgv_uebung_5/ApplicationWindow.ui b/hscgv_uebung_5/ApplicationWindow.ui index d67c5a3..038d964 100644 --- a/hscgv_uebung_5/ApplicationWindow.ui +++ b/hscgv_uebung_5/ApplicationWindow.ui @@ -12,7 +12,7 @@ - Aufgabe 1: "Ui, OpenGL!" + Aufgabe 5: "Jeros CUDA Realtime Raytracer!" @@ -21,7 +21,7 @@ 0 0 520 - 22 + 18 From b8019eb75dc63f8708b2dc1bdb276205b57f60d0 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 22 May 2015 18:05:00 +0200 Subject: [PATCH 24/87] update ctors --- hscgv_uebung_5/ray.h | 3 ++- hscgv_uebung_5/ray.inl | 12 ++++++++++++ 2 files changed, 14 insertions(+), 1 deletion(-) diff --git a/hscgv_uebung_5/ray.h b/hscgv_uebung_5/ray.h index e49e304..c8211df 100644 --- a/hscgv_uebung_5/ray.h +++ b/hscgv_uebung_5/ray.h @@ -25,7 +25,8 @@ class Ray // CONSTRUCTORS __host__ __device__ Ray(); __host__ __device__ Ray(const Ray &r); - __host__ __device__ Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll); + __host__ Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll); + __device__ Ray(const Vec3d o, const Vec3d d, unsigned int i, GeoObject* ol, LightObject* ll); __host__ __device__ ~Ray(); const Color __host__ shade() const; diff --git a/hscgv_uebung_5/ray.inl b/hscgv_uebung_5/ray.inl index c3d47ed..08a17ec 100644 --- a/hscgv_uebung_5/ray.inl +++ b/hscgv_uebung_5/ray.inl @@ -46,6 +46,18 @@ inline Ray::Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector Date: Fri, 22 May 2015 18:05:45 +0200 Subject: [PATCH 25/87] implement cuda specific shade method --- hscgv_uebung_5/ray.h | 2 +- hscgv_uebung_5/ray.inl | 28 +++++++++++++--------------- 2 files changed, 14 insertions(+), 16 deletions(-) diff --git a/hscgv_uebung_5/ray.h b/hscgv_uebung_5/ray.h index c8211df..421a468 100644 --- a/hscgv_uebung_5/ray.h +++ b/hscgv_uebung_5/ray.h @@ -30,7 +30,7 @@ class Ray __host__ __device__ ~Ray(); const Color __host__ shade() const; - const Color __device__ cudaShade(GeoObject* m_objListCuda, int objListSize, LightObject *m_lightListCuda, int lightListSize) const; + const Color __device__ shade(Ray *thisRay, Vec3d d_origin, Vec3d d_direction, GeoObject *d_objList, int objListSize, LightObject *d_lightList, int lightListSize, Color background) ; // access methods Vec3d __host__ __device__ origin() const; diff --git a/hscgv_uebung_5/ray.inl b/hscgv_uebung_5/ray.inl index 08a17ec..9bb5662 100644 --- a/hscgv_uebung_5/ray.inl +++ b/hscgv_uebung_5/ray.inl @@ -137,10 +137,8 @@ Ray::shade() const return Color(currentColor); } } -// Description: -// Determine color of this ray by tracing through the scene -inline const Color __device__ -Ray::cudaShade(GeoObject* m_objListCuda, int objListSize, LightObject* m_lightListCuda, int lightListSize) const + +inline const Color __device__ Ray::shade(Ray *thisRay, Vec3d d_origin, Vec3d d_direction, GeoObject *d_objList, int objListSize, LightObject *d_lightList, int lightListSize, Color background) { Color currentColor(0.0); for (int i=0; i<5; i++) { @@ -150,17 +148,17 @@ Ray::cudaShade(GeoObject* m_objListCuda, int objListSize, LightObject* m_lightLi // find closest object that intersects for (int j=0; j ray goes to infinity if (closest == NULL) { - if (m_depth == 0) { -// return g_sceneCuda.picture.background; // background color + if (i == 0) { + return background; // background color } else { return Color(0.0); // black @@ -168,23 +166,23 @@ Ray::cudaShade(GeoObject* m_objListCuda, int objListSize, LightObject* m_lightLi } else { // reflection - Vec3d intersectionPosition(m_origin + (m_direction * tMin)); + Vec3d intersectionPosition(d_origin + (d_direction * tMin)); Vec3d normal(closest->getNormal(intersectionPosition)); Ray reflectedRay(intersectionPosition, - m_direction.getReflectedAt(normal).getNormalized(), - m_depth+1,*m_objList,*m_lightList); + d_direction.getReflectedAt(normal).getNormalized(), + i+1,d_objList,d_lightList); // calculate lighting for (int j=0; j 0.0) { something_intersected = true; break; @@ -194,7 +192,7 @@ Ray::cudaShade(GeoObject* m_objListCuda, int objListSize, LightObject* m_lightLi // is it visible ? if (! something_intersected) - currentColor += shadedColor(&m_lightListCuda[j], reflectedRay, normal, closest); + currentColor += shadedColor(&d_lightList[j], reflectedRay, normal, closest); } // for all lights From 357a15e8d86df3b31e2f588358ae5f96ec6e9ce5 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 22 May 2015 18:07:41 +0200 Subject: [PATCH 26/87] remove const attribute for device specific handle --- hscgv_uebung_5/ray.h | 3 --- hscgv_uebung_5/ray.inl | 51 +++++++++++++++++++++--------------------- 2 files changed, 26 insertions(+), 28 deletions(-) diff --git a/hscgv_uebung_5/ray.h b/hscgv_uebung_5/ray.h index 421a468..453413b 100644 --- a/hscgv_uebung_5/ray.h +++ b/hscgv_uebung_5/ray.h @@ -36,9 +36,6 @@ class Ray Vec3d __host__ __device__ origin() const; Vec3d __host__ __device__ direction() const; - protected: - const Color __host__ __device__ shadedColor(LightObject *light, const Ray &reflectedray, const Vec3d &normal, GeoObject *obj) const; - private: Vec3d m_origin; Vec3d m_direction; diff --git a/hscgv_uebung_5/ray.inl b/hscgv_uebung_5/ray.inl index 9bb5662..778b731 100644 --- a/hscgv_uebung_5/ray.inl +++ b/hscgv_uebung_5/ray.inl @@ -64,6 +64,32 @@ inline Ray::~Ray( ) { } +// Description: +// Determine color contribution of a lightsource +inline Color __device__ __host__ +shadedColor(LightObject *light, const Ray &reflectedRay, const Vec3d &normal, GeoObject *obj) +{ + double ldot = light->direction() | normal; + Color reflectedColor = Color(0.0); + + // lambertian reflection model + if (ldot > 0.0) + reflectedColor += obj->reflectance() * (light->color() * ldot); + + // updated with ambient lightning as in: + // [GENERALISED AMBIENT REFLECTION MODELS FOR LAMBERTIAN AND PHONG SURFACES, Xiaozheng Zhang and Yongsheng Gao] +// reflectedColor += obj->ambient() * g_sceneCuda.ambience; + + // specular part + double spec = reflectedRay.direction() | light->direction(); + if (spec > 0.0) { + spec = obj->specular() * pow(spec, obj->specularExp()); + reflectedColor += light->color() * spec; + } + + return Color(reflectedColor); +} + // Description: // Determine color of this ray by tracing through the scene inline const Color __host__ @@ -203,31 +229,6 @@ inline const Color __device__ Ray::shade(Ray *thisRay, Vec3d d_origin, Vec3d d_d return Color(currentColor); } -// Description: -// Determine color contribution of a lightsource -inline const Color __device__ __host__ -Ray::shadedColor(LightObject *light, const Ray &reflectedRay, const Vec3d &normal, GeoObject *obj) const -{ - double ldot = light->direction() | normal; - Color reflectedColor = Color(0.0); - - // lambertian reflection model - if (ldot > 0.0) - reflectedColor += obj->reflectance() * (light->color() * ldot); - - // updated with ambient lightning as in: - // [GENERALISED AMBIENT REFLECTION MODELS FOR LAMBERTIAN AND PHONG SURFACES, Xiaozheng Zhang and Yongsheng Gao] -// reflectedColor += obj->ambient() * g_sceneCuda.ambience; - - // specular part - double spec = reflectedRay.direction() | light->direction(); - if (spec > 0.0) { - spec = obj->specular() * pow(spec, obj->specularExp()); - reflectedColor += light->color() * spec; - } - - return Color(reflectedColor); -} inline Vec3d __device__ __host__ Ray::origin() const From 4998c7e6c112e5c5b0e444d0cd37a5b3cceefedb Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 22 May 2015 18:08:26 +0200 Subject: [PATCH 27/87] explicit host function --- hscgv_uebung_5/raytracer.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 2516c63..0edd08e 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -42,7 +42,7 @@ Raytracer::~Raytracer() { cleanUp(); } -__host__ void +void Raytracer::start(float *renderedScene, int xRes, int yRes) { // setup viewport, its origin is bottom left // setup camera coordsys From 7e0f6e9a294ba3543033d682e44b3e7274686c47 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 22 May 2015 18:08:33 +0200 Subject: [PATCH 28/87] remove testing --- hscgv_uebung_5/raytracer.cu | 3 --- 1 file changed, 3 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 0edd08e..653ee5e 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -62,10 +62,7 @@ Raytracer::start(float *renderedScene, int xRes, int yRes) { Vec3d bottomLeft = g_scene.view.eyepoint + eye_dir - deltaX*xRes/2 - deltaY*yRes/2; // normal ray tracing: the color of the center of a pixel is computed -#ifndef __CUDA_ARCH__ - fprintf(stderr, "%d threads on board \n", omp_get_num_threads()); #pragma omp parallel for schedule(dynamic) collapse(2) -#endif for (int sy=yRes ; sy > 0 ; --sy) { for (int sx=0 ; sx < xRes ; ++sx) { // the center of the pixel we are looking at right now From a2e353d6cd46d3b5f2f45b604275c05ed663a47d Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 22 May 2015 18:09:48 +0200 Subject: [PATCH 29/87] substitute raytracing host routine by two different ones --- hscgv_uebung_5/raytracer.cu | 135 +++++++++++++++++++++--------------- 1 file changed, 78 insertions(+), 57 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 653ee5e..7900613 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -110,71 +110,92 @@ Raytracer::start(float *renderedScene, int xRes, int yRes) { } // foreach y } - void -Raytracer::startCuda(float *renderedScene, int xRes, int yRes) { - // setup viewport, its origin is bottom left - // setup camera coordsys - Vec3d eye_dir = (g_scene.view.lookat - g_scene.view.eyepoint).getNormalized(); - Vec3d eye_right = (eye_dir^g_scene.view.up).getNormalized(); - Vec3d eye_up = eye_dir^eye_right*-1; - - // calculatehe dimensions of the viewport using the scene's camera - float height = 2 * tan(M_PI/180 * .5 * g_scene.view.fovy); - float width = height * g_scene.view.aspect; - - // compute delta steps in each direction - Vec3d deltaX = eye_right * (width / xRes); - Vec3d deltaY = eye_up * (height / yRes); - - // this should be bottom left - Vec3d bottomLeft = g_scene.view.eyepoint + eye_dir - deltaX*xRes/2 - deltaY*yRes/2; +Raytracer::startCuda(float *renderedScene, int xRes, int yRes) +{ + // create empty space for scene on gpu - // normal ray tracing: the color of the center of a pixel is computed - for (int sy=yRes ; sy > 0 ; --sy) { - for (int sx=0 ; sx < xRes ; ++sx) { - // the center of the pixel we are looking at right now - Vec3d point = bottomLeft + deltaX*sx + deltaY*sy + deltaX/2 + deltaY/2; + // updload current position of camera and the like, objects/lights/bgColor (should be done only once) - // the direction of our look - Vec3d dir = point - g_scene.view.eyepoint; + // RAY TRACING: + dim3 block(32, 16, 1); + dim3 grid(xRes/ block.x, yRes / block.y, 1); +// startCudaKernel<<>>block(); +//startCudaKernel(float *renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, bool d_antialiasing, Color backgroundCol) +} - // create ray from view.eyepoint to view.lookat - Ray theRay(g_scene.view.eyepoint,dir.getNormalized(),0,g_objectList,g_lightList); - // compute the color - Color col = theRay.shade(); - // in case we are using antialiasing, calculate the color of this pixel by averaging - if (m_antialiasing) { - // scale the midpoint color since we are going to use 5 points to average our color - col *= 0.2; +void __global__ +startCudaKernel(float *renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, bool d_antialiasing, Color backgroundCol) +{ + // find out id of this thread + unsigned int idx = blockIdx.x*blockDim.x + threadIdx.x; + if (idx>=xRes*yRes) + return; + // we are handling pixel (x,y) now + int sx = idx % xRes; + int sy = idx / xRes; + + // setup viewport, its origin is bottom left + // setup camera coordsys + Vec3d eye_dir = (lookat - eyepoint).getNormalized(); + Vec3d eye_right = (eye_dir^up).getNormalized(); + Vec3d eye_up = eye_dir^eye_right*-1; - // besides taking shooting a ray through the midpoint 'o', we calculate - // the pixels color by shooting 4 more rays through the points 'x' - // and averaging their values - // ------- - // | x x | - // | o | - // | x x | - // -------- - for(float dx = -1/4.; dx <= 1/4.; dx+=1/2.) { - for(float dy = -1/4.; dy <= 1/4.; dy+=1/2.) { - Vec3d superSamplePoint = point + deltaX*dx + deltaY*dy; - Vec3d superSampleDir = superSamplePoint - g_scene.view.eyepoint; + // calculatehe dimensions of the viewport using the scene's camera + float height = 2 * tan(M_PI/180 * .5 * fovy); + float width = height * aspect; - // create ray from view.eyepoint to view.lookat - Ray theRay(g_scene.view.eyepoint,superSampleDir.getNormalized(),0,g_objectList,g_lightList); - col += theRay.shade()*0.2;//color, recursive_ray_trace(eye, ray, 0)); - } - } - } + // compute delta steps in each direction + Vec3d deltaX = eye_right * (width / xRes); + Vec3d deltaY = eye_up * (height / yRes); - int index = 3*((sy-1) * xRes + sx); - renderedScene[index + 0] = col[0]; - renderedScene[index + 1] = col[1]; - renderedScene[index + 2] = col[2]; - } // foreach x - } // foreach y + // this should be bottom left + Vec3d bottomLeft = eyepoint + eye_dir - deltaX*xRes/2 - deltaY*yRes/2; + + // the center of the pixel we are looking at right now + Vec3d point = bottomLeft + deltaX*sx + deltaY*sy + deltaX/2 + deltaY/2; + + // the direction of our look + Vec3d dir = point - eyepoint; + + // create ray from view.eyepoint to view.lookat + Ray theRay(eyepoint,dir.getNormalized(),0,d_objectList,d_lightList); + + // compute the color + Color col; + theRay.shade(&theRay, eyepoint, point, d_objectList, objListSize, d_lightList, lightListSize, backgroundCol); + + // in case we are using antialiasing, calculate the color of this pixel by averaging + if (d_antialiasing) { + // scale the midpoint color since we are going to use 5 points to average our color + col *= 0.2; + + // besides taking shooting a ray through the midpoint 'o', we calculate + // the pixels color by shooting 4 more rays through the points 'x' + // and averaging their values + // ------- + // | x x | + // | o | + // | x x | + // -------- + for(float dx = -1/4.; dx <= 1/4.; dx+=1/2.) { + for(float dy = -1/4.; dy <= 1/4.; dy+=1/2.) { + Vec3d superSamplePoint = point + deltaX*dx + deltaY*dy; + Vec3d superSampleDir = superSamplePoint - eyepoint; + + // create ray from view.eyepoint to view.lookat + Ray theRay(eyepoint,superSampleDir.getNormalized(),0,d_objectList,d_lightList); + col += theRay.shade(&theRay,eyepoint,superSamplePoint, d_objectList, objListSize, d_lightList, lightListSize, backgroundCol)*0.2; + //color, recursive_ray_trace(eye, ray, 0)); + } + } + } + + int index = 3*((sy-1) * xRes + sx); + renderedScene[index + 0] = col[0]; + renderedScene[index + 1] = col[1]; + renderedScene[index + 2] = col[2]; } From d677befd12cf1d245310f0d7814aa2bf9b649ed8 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 22 May 2015 18:09:58 +0200 Subject: [PATCH 30/87] add gpu constants --- hscgv_uebung_5/raytracer.cu | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 7900613..7171df9 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -124,6 +124,12 @@ Raytracer::startCuda(float *renderedScene, int xRes, int yRes) //startCudaKernel(float *renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, bool d_antialiasing, Color backgroundCol) } +// we can save constants on the GPU in an extra space with a lot faster access +__constant__ GeoObject* d_objectList; +__constant__ LightObject* d_lightList; +__constant__ int objListSize; +__constant__ int lightListSize; +__constant__ double fovy; void __global__ From ed6efc35dfe7a07e914e38559f50654471013f9f Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Fri, 22 May 2015 18:10:13 +0200 Subject: [PATCH 31/87] update aspect vector on resize --- hscgv_uebung_5/GLFrame.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 2e770b1..5092146 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -104,6 +104,7 @@ void GLFrame::loadTexture() void GLFrame::resizeGL(int w, int h) { m_aspect = (GLdouble)w / (GLdouble)h; + g_scene.view.aspect = m_aspect; m_width = w; m_height = h; From 4b9014d213f7ee0080430865fa16debce5efe73e Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 26 May 2015 11:33:23 +0200 Subject: [PATCH 32/87] remove unused counters --- hscgv_uebung_5/input.y | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/hscgv_uebung_5/input.y b/hscgv_uebung_5/input.y index 3f08614..88a9eb7 100644 --- a/hscgv_uebung_5/input.y +++ b/hscgv_uebung_5/input.y @@ -341,8 +341,6 @@ static int surfaceCounter = 0; static int propertyCounter = 0; static int objectCounter = 0; static int lightCounter = 0; -static int vertexCounter = 0; -static int indexCounter = 0; /* * checking to make sure indices are okay @@ -351,8 +349,7 @@ static int nproperties = 0; static int nsurfaces = 0; static int nobjs = 0; static int nlights = 0; -static int nvertices = 0; -static int nindices = 0; + /* * so that we can keep track if a particular parameter appeared From a228e65c6796ad4af2e95af8c3dde0c7e8a52d04 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 26 May 2015 11:40:55 +0200 Subject: [PATCH 33/87] remove explicit initialization of object and light lists --- hscgv_uebung_5/ray.h | 2 +- hscgv_uebung_5/ray.inl | 8 +++----- hscgv_uebung_5/raytracer.cu | 8 ++++---- 3 files changed, 8 insertions(+), 10 deletions(-) diff --git a/hscgv_uebung_5/ray.h b/hscgv_uebung_5/ray.h index 453413b..bd9a5fc 100644 --- a/hscgv_uebung_5/ray.h +++ b/hscgv_uebung_5/ray.h @@ -26,7 +26,7 @@ class Ray __host__ __device__ Ray(); __host__ __device__ Ray(const Ray &r); __host__ Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll); - __device__ Ray(const Vec3d o, const Vec3d d, unsigned int i, GeoObject* ol, LightObject* ll); + __device__ Ray(const Vec3d o, const Vec3d d, unsigned int i); __host__ __device__ ~Ray(); const Color __host__ shade() const; diff --git a/hscgv_uebung_5/ray.inl b/hscgv_uebung_5/ray.inl index 778b731..a829557 100644 --- a/hscgv_uebung_5/ray.inl +++ b/hscgv_uebung_5/ray.inl @@ -48,14 +48,12 @@ inline Ray::Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vectorgetNormal(intersectionPosition)); Ray reflectedRay(intersectionPosition, d_direction.getReflectedAt(normal).getNormalized(), - i+1,d_objList,d_lightList); + i+1); // calculate lighting for (int j=0; j Date: Tue, 26 May 2015 11:42:37 +0200 Subject: [PATCH 34/87] removed semicolons --- hscgv_uebung_5/geoobject.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/hscgv_uebung_5/geoobject.h b/hscgv_uebung_5/geoobject.h index 0e3ecd8..18cfa46 100644 --- a/hscgv_uebung_5/geoobject.h +++ b/hscgv_uebung_5/geoobject.h @@ -22,15 +22,15 @@ class GeoObjectProperties { public: // CONSTRUCTORS - GeoObjectProperties() {}; + GeoObjectProperties() {} GeoObjectProperties(const Vec3d& a,const Vec3d& r, double s, int e, double m) { m_ambient = a; m_reflectance = r; m_specular = s; m_specularExp = e; m_mirror = m; - }; - ~GeoObjectProperties() {}; + } + ~GeoObjectProperties() {} // access methods Color __host__ __device__ ambient() const; @@ -55,8 +55,8 @@ class GeoObject { public: // CONSTRUCTORS - GeoObject() {}; - virtual ~GeoObject() {}; + GeoObject() {} + virtual ~GeoObject() {} virtual Vec3d __host__ __device__ getNormal(const Vec3d &v) const = 0; virtual double __host__ __device__ intersect(const Ray &r) const = 0; From 91f759f910428fa585453b274818c3e13651805a Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 26 May 2015 11:43:14 +0200 Subject: [PATCH 35/87] move LightObjectProperties out of the class - as it is in GeoObject --- hscgv_uebung_5/lightobject.h | 13 +++++-------- 1 file changed, 5 insertions(+), 8 deletions(-) diff --git a/hscgv_uebung_5/lightobject.h b/hscgv_uebung_5/lightobject.h index 79ed37d..e9e6728 100644 --- a/hscgv_uebung_5/lightobject.h +++ b/hscgv_uebung_5/lightobject.h @@ -12,6 +12,11 @@ #include "vector.h" +typedef struct { + Color color; + Vec3d direction; +} LightObjectProperties; + //! A light source /*! This class describes light sources positioned at infinity. */ @@ -26,14 +31,6 @@ class LightObject // access methods Color __device__ __host__ color() const; Vec3d __device__ __host__ direction() const; - - protected: - typedef struct { - Color color; - Vec3d direction; - } LightObjectProperties; - - private: LightObjectProperties *m_properties; }; From ac09a9153fabb3018cc69b2503d1d4e7ab9523e6 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 26 May 2015 11:44:01 +0200 Subject: [PATCH 36/87] add cuda specific members --- hscgv_uebung_5/raytracer.h | 12 +++++++++++- 1 file changed, 11 insertions(+), 1 deletion(-) diff --git a/hscgv_uebung_5/raytracer.h b/hscgv_uebung_5/raytracer.h index 35f483a..223d67c 100644 --- a/hscgv_uebung_5/raytracer.h +++ b/hscgv_uebung_5/raytracer.h @@ -28,8 +28,18 @@ class Raytracer Raytracer(const char* filename, bool antialiasing); ~Raytracer(); void start(float* renderedScene, int xRes, int yRes); - void startCuda(float* renderedScene, int xRes, int yRes); + void startCuda(float *renderedScene, int xRes, int yRes); + void __device__ initData(); + void __host__ initCuda(); bool m_isFileLoaded; + GeoObject* d_objList; + GeoObjectProperties* d_objPropList; + int d_objListSize; + LightObject* d_lightList; + LightObjectProperties* d_lightPropList; + int d_lightListSize; + double d_fovy; + private: const char *m_filename; bool m_antialiasing; From d11ccdfb03af6064794a001ad16b5625130135d4 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 26 May 2015 11:52:08 +0200 Subject: [PATCH 37/87] remove constants, add initCuda and initData and update Kernel --- hscgv_uebung_5/raytracer.cu | 39 +++++++++++++++++++++++++++++-------- 1 file changed, 31 insertions(+), 8 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 979265a..12b1281 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -124,16 +124,39 @@ Raytracer::startCuda(float *renderedScene, int xRes, int yRes) //startCudaKernel(float *renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, bool d_antialiasing, Color backgroundCol) } -// we can save constants on the GPU in an extra space with a lot faster access -__constant__ GeoObject* d_objectList; -__constant__ LightObject* d_lightList; -__constant__ int objListSize; -__constant__ int lightListSize; -__constant__ double fovy; +#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } +inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true) { + if (code != cudaSuccess) { + fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); + if (abort) exit(code); + } +} +//! we need some kind of initialization of our device +void Raytracer::initCuda() { + // get some space for the objects and their properties ( + gpuErrchk (cudaMalloc((void **) &d_objList, sizeof(GeoObject) * g_objectList.size())); + gpuErrchk (cudaMalloc((void **) &d_objPropList, sizeof(GeoObjectProperties) * g_objectList.size())); + gpuErrchk (cudaMalloc((void **) &d_lightList, sizeof(LightObject) * g_lightList.size())); + gpuErrchk (cudaMalloc((void **) &d_lightPropList, sizeof(LightObjectProperties) * g_lightList.size())); + + gpuErrchk (cudaMemcpy(d_objList, g_objectList.data(), sizeof(GeoObject) * g_objectList.size(), cudaMemcpyHostToDevice)); + gpuErrchk (cudaMemcpy(d_objPropList, g_objectList.data(), sizeof(GeoObjectProperties) * g_objectList.size(), cudaMemcpyHostToDevice)); + gpuErrchk (cudaMemcpy(d_lightList, g_objectList.data(), sizeof(LightObject) * g_lightList.size(), cudaMemcpyHostToDevice)); + gpuErrchk (cudaMemcpy(d_lightPropList, g_objectList.data(), sizeof(LightObjectProperties) * g_lightList.size(), cudaMemcpyHostToDevice)); + } + +void __device__ +Raytracer::initData() { + for (int i=0; icolor = d_lightList[i].color(); + d_lightList[i].m_properties->direction = d_lightList[i].direction(); + } +} void __global__ -startCudaKernel(float *renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, bool d_antialiasing, Color backgroundCol) +startCudaKernel(float *renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, bool d_antialiasing, Color backgroundCol, GeoObject* d_objList, int d_objListSize, LightObject* d_lightList, int d_lightListSize, double d_fovy) { // find out id of this thread unsigned int idx = blockIdx.x*blockDim.x + threadIdx.x; @@ -150,7 +173,7 @@ startCudaKernel(float *renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d Vec3d eye_up = eye_dir^eye_right*-1; // calculatehe dimensions of the viewport using the scene's camera - float height = 2 * tan(M_PI/180 * .5 * fovy); + float height = 2 * tan(M_PI/180 * .5 * d_fovy); float width = height * aspect; // compute delta steps in each direction From 4bbebcb55f2d3c8d4b31bea7e2a52a269d8526bc Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 26 May 2015 16:26:54 +0200 Subject: [PATCH 38/87] ignore compiler errors due to host and device code mixup --- hscgv_uebung_5/types.h | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/hscgv_uebung_5/types.h b/hscgv_uebung_5/types.h index 7ced9a9..d7a1175 100644 --- a/hscgv_uebung_5/types.h +++ b/hscgv_uebung_5/types.h @@ -13,6 +13,12 @@ #include #include +// ignore all compiler errors due to mixing up of cuda host and device code +#ifndef __CUDACC__ +#define __host__ +#define __device__ +#endif + //! error message macro /*! Use this macro to generate informative output in case the program detects From 57e64461f5b60f40b0ea8b3f435655b1f8bbb9f1 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 26 May 2015 18:26:39 +0200 Subject: [PATCH 39/87] geiler commit --- hscgv_uebung_5/lightobject.h | 3 ++ hscgv_uebung_5/lightobject.inl | 7 +++ hscgv_uebung_5/raytracer.cu | 87 +++++++++++++++++----------------- hscgv_uebung_5/raytracer.h | 9 ++-- 4 files changed, 58 insertions(+), 48 deletions(-) diff --git a/hscgv_uebung_5/lightobject.h b/hscgv_uebung_5/lightobject.h index e9e6728..ba959fa 100644 --- a/hscgv_uebung_5/lightobject.h +++ b/hscgv_uebung_5/lightobject.h @@ -31,6 +31,9 @@ class LightObject // access methods Color __device__ __host__ color() const; Vec3d __device__ __host__ direction() const; + void __host__ __device__ setProperties(LightObjectProperties *p); + + protected: LightObjectProperties *m_properties; }; diff --git a/hscgv_uebung_5/lightobject.inl b/hscgv_uebung_5/lightobject.inl index 7327df4..d168807 100644 --- a/hscgv_uebung_5/lightobject.inl +++ b/hscgv_uebung_5/lightobject.inl @@ -48,3 +48,10 @@ LightObject::direction() const // std::cerr<<"WARNING: properties not set"<>>block(); -//startCudaKernel(float *renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, bool d_antialiasing, Color backgroundCol) -} - #define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true) { if (code != cudaSuccess) { @@ -132,31 +118,18 @@ inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=t } } -//! we need some kind of initialization of our device -void Raytracer::initCuda() { - // get some space for the objects and their properties ( - gpuErrchk (cudaMalloc((void **) &d_objList, sizeof(GeoObject) * g_objectList.size())); - gpuErrchk (cudaMalloc((void **) &d_objPropList, sizeof(GeoObjectProperties) * g_objectList.size())); - gpuErrchk (cudaMalloc((void **) &d_lightList, sizeof(LightObject) * g_lightList.size())); - gpuErrchk (cudaMalloc((void **) &d_lightPropList, sizeof(LightObjectProperties) * g_lightList.size())); - - gpuErrchk (cudaMemcpy(d_objList, g_objectList.data(), sizeof(GeoObject) * g_objectList.size(), cudaMemcpyHostToDevice)); - gpuErrchk (cudaMemcpy(d_objPropList, g_objectList.data(), sizeof(GeoObjectProperties) * g_objectList.size(), cudaMemcpyHostToDevice)); - gpuErrchk (cudaMemcpy(d_lightList, g_objectList.data(), sizeof(LightObject) * g_lightList.size(), cudaMemcpyHostToDevice)); - gpuErrchk (cudaMemcpy(d_lightPropList, g_objectList.data(), sizeof(LightObjectProperties) * g_lightList.size(), cudaMemcpyHostToDevice)); - } - -void __device__ -Raytracer::initData() { - for (int i=0; icolor = d_lightList[i].color(); - d_lightList[i].m_properties->direction = d_lightList[i].direction(); - } +void __global__ +initPropertiesKernel(GeoObject* d_objList, GeoObjectProperties* d_objPropList, int d_objListSize, LightObject* d_lightList, LightObjectProperties* d_lightPropList, int d_lightListSize) { + // setup the objects properties one by one, since we needed to copy them by hand into different lists (objPropList and lightPropList) + for (int i=0; i>>(d_objList, d_objPropList, g_objectList.size(), d_lightList, d_lightPropList, g_lightList.size()); + } + +void +Raytracer::startCuda(float *renderedScene, int xRes, int yRes) +{ + // RAY TRACING: + dim3 block(32, 16, 1); + dim3 grid(xRes/ block.x, yRes / block.y, 1); + startCudaKernel<<>>(d_renderedScene, xRes, yRes, + g_scene.view.eyepoint, g_scene.view.up, g_scene.view.lookat, g_scene.view.aspect, g_scene.view.fovy, g_scene.picture.background, + m_antialiasing, d_objList, g_objectList.size(), d_lightList, g_lightList.size()); +} diff --git a/hscgv_uebung_5/raytracer.h b/hscgv_uebung_5/raytracer.h index 223d67c..133864c 100644 --- a/hscgv_uebung_5/raytracer.h +++ b/hscgv_uebung_5/raytracer.h @@ -29,16 +29,15 @@ class Raytracer ~Raytracer(); void start(float* renderedScene, int xRes, int yRes); void startCuda(float *renderedScene, int xRes, int yRes); - void __device__ initData(); - void __host__ initCuda(); + void initCuda(); bool m_isFileLoaded; + + //! CUDA Properties GeoObject* d_objList; GeoObjectProperties* d_objPropList; - int d_objListSize; LightObject* d_lightList; LightObjectProperties* d_lightPropList; - int d_lightListSize; - double d_fovy; + float *d_renderedScene; private: const char *m_filename; From da5ea49cbe8f5c575130d3c05619f5d878e942e2 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 26 May 2015 21:12:30 +0200 Subject: [PATCH 40/87] setup compilation for mac --- hscgv_uebung_5/cuda.pri | 12 +++++++++--- hscgv_uebung_5/raytracer.cu | 2 ++ hscgv_uebung_5/src.pro | 6 +++--- 3 files changed, 14 insertions(+), 6 deletions(-) diff --git a/hscgv_uebung_5/cuda.pri b/hscgv_uebung_5/cuda.pri index 0b653ac..f3faa71 100644 --- a/hscgv_uebung_5/cuda.pri +++ b/hscgv_uebung_5/cuda.pri @@ -11,12 +11,18 @@ win32 { cuda.commands = $(CUDA_BIN_DIR)/nvcc.exe -c -Xcompiler $$join(QMAKE_CXXFLAGS,",") $$join(INCLUDEPATH,'" -I "','-I "','"') ${QMAKE_FILE_NAME} -o ${QMAKE_FILE_OUT} } + +unix:!macx:{ + CUDA_DIR = /usr/local/cuda-5.5 + QMAKE_LIBDIR += $$CUDA_DIR/lib64 +} +macx:{ + CUDA_DIR = /Developer/NVIDIA/CUDA-7.0 + QMAKE_LIBDIR += $$CUDA_DIR/lib +} unix { # auto-detect CUDA path - CUDA_DIR = $$system(which nvcc | sed 's,/bin/nvcc$,,') - CUDA_DIR = /usr/local/cuda-5.5 INCLUDEPATH += $$CUDA_DIR/include - QMAKE_LIBDIR += $$CUDA_DIR/lib64 LIBS += -lcudart NVCCFLAGS="-use_fast_math --ptxas-options=-v" NVCCFLAGS+="-Xptxas -fastimul -arch=sm_20" # 24 bit integer multiplication should be sufficient diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 49d558e..f2c3a97 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -5,7 +5,9 @@ */ #include "raytracer.h" +#ifndef __APPLE__ #include "omp.h" +#endif //! get scene description from the parser /*! diff --git a/hscgv_uebung_5/src.pro b/hscgv_uebung_5/src.pro index 695fdd0..685996e 100644 --- a/hscgv_uebung_5/src.pro +++ b/hscgv_uebung_5/src.pro @@ -25,16 +25,16 @@ DEFINES *= TRACE VERBOSE CONFIG += debug CONFIG -= release -QMAKE_CXXFLAGS += -W -Wall -fopenmp -QMAKE_LFLAGS += -fopenmp FORMS += \ ApplicationWindow.ui unix:!macx:{ LIBS *= -lGLEW -lglut -lGLU -include(cuda.pri) +QMAKE_CXXFLAGS += -W -Wall -fopenmp +QMAKE_LFLAGS += -fopenmp } +include(cuda.pri) macx { LIBS *= -lGLEW -L/usr/local/opt/glew/lib From 14d94b123ecc4d0b57872fed375823446479366c Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 26 May 2015 21:13:08 +0200 Subject: [PATCH 41/87] change naming to be more useful --- hscgv_uebung_5/raytracer.cu | 12 +++++++----- hscgv_uebung_5/raytracer.h | 4 ++-- 2 files changed, 9 insertions(+), 7 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index f2c3a97..f386e8e 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -45,7 +45,7 @@ Raytracer::~Raytracer() { } void -Raytracer::start(float *renderedScene, int xRes, int yRes) { +Raytracer::render(float *renderedScene, int xRes, int yRes) { // setup viewport, its origin is bottom left // setup camera coordsys Vec3d eye_dir = (g_scene.view.lookat - g_scene.view.eyepoint).getNormalized(); @@ -130,7 +130,7 @@ initPropertiesKernel(GeoObject* d_objList, GeoObjectProperties* d_objPropList, i } void __global__ -startCudaKernel(float *d_renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, double fovy, Color backgroundCol, bool antialiasing, +renderKernel(float *d_renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, double fovy, Color backgroundCol, bool antialiasing, GeoObject* d_objList, int objListSize, LightObject* d_lightList, int lightListSize) { // find out id of this thread @@ -206,7 +206,8 @@ startCudaKernel(float *d_renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3 //! we need some kind of initialization of our device -void Raytracer::initCuda() { +void +Raytracer::initCuda() { // get some space for the objects and their properties ( gpuErrchk (cudaMalloc((void **) &d_objList, sizeof(GeoObject) * g_objectList.size())); gpuErrchk (cudaMalloc((void **) &d_objPropList, sizeof(GeoObjectProperties) * g_objectList.size())); @@ -221,13 +222,14 @@ void Raytracer::initCuda() { initPropertiesKernel<<<1,1>>>(d_objList, d_objPropList, g_objectList.size(), d_lightList, d_lightPropList, g_lightList.size()); } +//! start the rendering routine on the device void -Raytracer::startCuda(float *renderedScene, int xRes, int yRes) +Raytracer::renderCuda(float *renderedScene, int xRes, int yRes) { // RAY TRACING: dim3 block(32, 16, 1); dim3 grid(xRes/ block.x, yRes / block.y, 1); - startCudaKernel<<>>(d_renderedScene, xRes, yRes, + renderKernel<<>>(d_renderedScene, xRes, yRes, g_scene.view.eyepoint, g_scene.view.up, g_scene.view.lookat, g_scene.view.aspect, g_scene.view.fovy, g_scene.picture.background, m_antialiasing, d_objList, g_objectList.size(), d_lightList, g_lightList.size()); } diff --git a/hscgv_uebung_5/raytracer.h b/hscgv_uebung_5/raytracer.h index 133864c..9f70ded 100644 --- a/hscgv_uebung_5/raytracer.h +++ b/hscgv_uebung_5/raytracer.h @@ -27,8 +27,8 @@ class Raytracer Raytracer(bool antialiasing); Raytracer(const char* filename, bool antialiasing); ~Raytracer(); - void start(float* renderedScene, int xRes, int yRes); - void startCuda(float *renderedScene, int xRes, int yRes); + void render(float* renderedScene, int xRes, int yRes); + void renderCuda(float *renderedScene, int xRes, int yRes); void initCuda(); bool m_isFileLoaded; From a7942d85c560bb3721844c865d297254a387d4c8 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 26 May 2015 21:14:00 +0200 Subject: [PATCH 42/87] introduce PixelBufferObject drawing on GPU --- hscgv_uebung_5/GLFrame.cpp | 54 ++++++++++++++++++++++++++++++++++---- hscgv_uebung_5/GLFrame.h | 5 ++++ 2 files changed, 54 insertions(+), 5 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 5092146..6e2a112 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -14,6 +14,9 @@ #include #include +#include +#include +#include #include "GLFrame.h" @@ -89,12 +92,17 @@ void GLFrame::loadTexture() glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER); + // move data into a texture - the last parameter is either host data or + // NULL since we only want to allocate memory, not initialize it in case of PBO +// glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, m_width, m_height, 0, GL_RGB, GL_FLOAT, +// (m_renderMode==CPU)?((GLvoid*)m_data):NULL); CHECKGL; // move data into a texture - glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, m_width, m_height, 0, GL_RGB, GL_FLOAT, (GLvoid*)m_data);CHECKGL; + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, m_width, m_height, 0, GL_RGB, GL_FLOAT, (GLvoid*)m_data); CHECKGL; // free data since its on the texture free(m_data); +// cudaGLMapBufferObject((void**)&m_currentDevice, m_currentVBO_CUDA); glBindTexture(GL_TEXTURE_2D, 0); } @@ -428,10 +436,21 @@ void GLFrame::setRenderMode(int mode) m_renderMode = (RenderMode)mode; if(mode == CPU) { - + // unset VBO + cudaGraphicsUnregisterResource(m_currentVBO_CUDA); + glDeleteBuffers(1, &m_currentVBO); } else if (mode == GPU) { - // Todo: setup cuda (maybe earlier) and push all data to GPU + // setup VBO + cudaGLSetGLDevice(0); + // Create buffer object and register it with CUDA + glGenBuffers(1, &m_currentVBO); + glBindBuffer(GL_ARRAY_BUFFER, m_currentVBO); + unsigned int size = m_width * m_height * 3;//? * sizeof(float); + glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW); + glBindBuffer(GL_ARRAY_BUFFER, 0); + cudaGraphicsGLRegisterBuffer(&m_currentVBO_CUDA, m_currentVBO, cudaGraphicsMapFlagsWriteDiscard); + } m_raytracingNeeded = true; @@ -558,11 +577,32 @@ void GLFrame::renderScene() { case CPU: m_data = (float*)malloc(sizeof(float) * 3 * m_width * m_height); - m_raytracer->start(m_data, m_width, m_height); - m_raytracingNeeded = false; + m_raytracer->render(m_data, m_width, m_height); loadTexture(); + m_raytracingNeeded = false; break; case GPU: + // Map buffer object for writing from CUDA + cudaGraphicsMapResources(1, &m_currentVBO_CUDA, 0); + size_t num_bytes; + cudaGraphicsResourceGetMappedPointer((void**)&m_data, + &num_bytes, + m_currentVBO_CUDA); + // execute rendering on gpu with kernel + m_raytracer->renderCuda(m_data, m_width, m_height); + // Unmap buffer object + cudaGraphicsUnmapResources(1, &m_currentVBO_CUDA, 0); + // Render from buffer object + glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); + glBindBuffer(GL_ARRAY_BUFFER, m_currentVBO); + glVertexPointer(3, GL_FLOAT, 0, 0); + glEnableClientState(GL_VERTEX_ARRAY); + glDrawArrays(GL_POINTS, 0, m_width * m_height); + glDisableClientState(GL_VERTEX_ARRAY); + // Swap buffers + swapBuffers(); + // indicate done job + m_raytracingNeeded = false; break; } } @@ -570,8 +610,12 @@ void GLFrame::renderScene() void GLFrame::loadScene(const QString &filename) { + // create a new raytracer delete m_raytracer; m_raytracer = new Raytracer(filename.toStdString().c_str(), m_antialiasing); + // initialize the cuda structure + m_raytracer->initCuda(); + // and beg for drawing m_raytracingNeeded = true; updateGL(); } diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index 7a484d9..03ae08b 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -176,6 +176,11 @@ public slots: //! the currently used texture cache GLuint m_texHandle; + //! the currently used VBO + GLuint m_currentVBO; + //! the currently used VBO in CUDA + struct cudaGraphicsResource* m_currentVBO_CUDA; + //! axes visibility bool m_axesVisible; //! model visibility From a13de08e822e102479c8f74925e77cdffe0ee60e Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:17:13 +0200 Subject: [PATCH 43/87] call it cudadata --- hscgv_uebung_5/GLFrame.h | 1 + hscgv_uebung_5/raytracer.h | 3 +-- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index 03ae08b..8eafa17 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -193,6 +193,7 @@ public slots: Raytracer *m_raytracer; bool m_raytracingNeeded; float* m_data; + float* m_cudaData; //! how our model is to be drawn RenderMode m_renderMode; diff --git a/hscgv_uebung_5/raytracer.h b/hscgv_uebung_5/raytracer.h index 9f70ded..67acd04 100644 --- a/hscgv_uebung_5/raytracer.h +++ b/hscgv_uebung_5/raytracer.h @@ -28,7 +28,7 @@ class Raytracer Raytracer(const char* filename, bool antialiasing); ~Raytracer(); void render(float* renderedScene, int xRes, int yRes); - void renderCuda(float *renderedScene, int xRes, int yRes); + void renderCuda(float *cudaData, int xRes, int yRes); void initCuda(); bool m_isFileLoaded; @@ -37,7 +37,6 @@ class Raytracer GeoObjectProperties* d_objPropList; LightObject* d_lightList; LightObjectProperties* d_lightPropList; - float *d_renderedScene; private: const char *m_filename; From ca8ff343d8b2009eb959c0dd57c433d11ebe83ab Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:17:42 +0200 Subject: [PATCH 44/87] free data structures on exit --- hscgv_uebung_5/GLFrame.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 6e2a112..8a41270 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -58,6 +58,9 @@ GLFrame::GLFrame(ApplicationWindow *parent) GLFrame::~GLFrame() { glDeleteTextures(1, &m_texHandle); + free(m_data); + cudaFree(m_cudaData); + cudaDeviceReset(); } // ---------------------------- Basic OpenGL Widget Methods ------------------ From d07ae6b357edbe5d1601e83a6790f25853701a4b Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:19:56 +0200 Subject: [PATCH 45/87] update load texture --- hscgv_uebung_5/GLFrame.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 8a41270..aa09032 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -95,13 +95,12 @@ void GLFrame::loadTexture() glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER); - // move data into a texture - the last parameter is either host data or - // NULL since we only want to allocate memory, not initialize it in case of PBO -// glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, m_width, m_height, 0, GL_RGB, GL_FLOAT, -// (m_renderMode==CPU)?((GLvoid*)m_data):NULL); CHECKGL; - // move data into a texture glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, m_width, m_height, 0, GL_RGB, GL_FLOAT, (GLvoid*)m_data); CHECKGL; + if (m_renderMode == GPU) { + // register this texture with CUDA + } + // free data since its on the texture free(m_data); From ebcd2227be6d5e3e7d8c43b643b4a3597d9ee0b0 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:22:09 +0200 Subject: [PATCH 46/87] remove VBO stuff for CUDA --- hscgv_uebung_5/GLFrame.cpp | 33 --------------------------------- 1 file changed, 33 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index aa09032..6699aa1 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -439,20 +439,8 @@ void GLFrame::setRenderMode(int mode) if(mode == CPU) { // unset VBO - cudaGraphicsUnregisterResource(m_currentVBO_CUDA); - glDeleteBuffers(1, &m_currentVBO); } else if (mode == GPU) { - // setup VBO - cudaGLSetGLDevice(0); - // Create buffer object and register it with CUDA - glGenBuffers(1, &m_currentVBO); - glBindBuffer(GL_ARRAY_BUFFER, m_currentVBO); - unsigned int size = m_width * m_height * 3;//? * sizeof(float); - glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW); - glBindBuffer(GL_ARRAY_BUFFER, 0); - cudaGraphicsGLRegisterBuffer(&m_currentVBO_CUDA, m_currentVBO, cudaGraphicsMapFlagsWriteDiscard); - } m_raytracingNeeded = true; @@ -584,27 +572,6 @@ void GLFrame::renderScene() m_raytracingNeeded = false; break; case GPU: - // Map buffer object for writing from CUDA - cudaGraphicsMapResources(1, &m_currentVBO_CUDA, 0); - size_t num_bytes; - cudaGraphicsResourceGetMappedPointer((void**)&m_data, - &num_bytes, - m_currentVBO_CUDA); - // execute rendering on gpu with kernel - m_raytracer->renderCuda(m_data, m_width, m_height); - // Unmap buffer object - cudaGraphicsUnmapResources(1, &m_currentVBO_CUDA, 0); - // Render from buffer object - glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); - glBindBuffer(GL_ARRAY_BUFFER, m_currentVBO); - glVertexPointer(3, GL_FLOAT, 0, 0); - glEnableClientState(GL_VERTEX_ARRAY); - glDrawArrays(GL_POINTS, 0, m_width * m_height); - glDisableClientState(GL_VERTEX_ARRAY); - // Swap buffers - swapBuffers(); - // indicate done job - m_raytracingNeeded = false; break; } } From 4adc22f2271786e34cfe01e01cd51652f598dc3b Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:22:24 +0200 Subject: [PATCH 47/87] we use floating point textures --- hscgv_uebung_5/GLFrame.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 6699aa1..6ce101e 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -537,7 +537,7 @@ void GLFrame::drawFullScreenQuad() glBindTexture(GL_TEXTURE_2D, m_texHandle); - glTexEnvi(GL_TEXTURE_ENV, GL_TEXTURE_ENV_MODE, GL_REPLACE); + glTexEnvf(GL_TEXTURE_ENV, GL_TEXTURE_ENV_MODE, GL_REPLACE); // set to red glColor3f(1.0,0.0,.0); From 361cbf2b670e2f4fa6581d8fc8613cf6ee4671ed Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:26:03 +0200 Subject: [PATCH 48/87] handle raytracing via copying data to GPU and back --- hscgv_uebung_5/GLFrame.cpp | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 6ce101e..8830cea 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -561,19 +561,25 @@ void GLFrame::renderScene() // only render the scene if there is some intialized raytracer if(!m_raytracingNeeded || !m_raytracer || !m_raytracer->m_isFileLoaded) return; - // draw scene + int sizeTex = sizeof(float) * 3 * m_width * m_height; + m_data = (float*)malloc(sizeTex); switch(m_renderMode) { case CPU: - m_data = (float*)malloc(sizeof(float) * 3 * m_width * m_height); m_raytracer->render(m_data, m_width, m_height); - loadTexture(); - m_raytracingNeeded = false; break; case GPU: + cudaMalloc(&m_cudaData,sizeTex); + m_raytracer->renderCuda(m_cudaData, m_width, m_height); + cudaGetLastError(); + cudaMemcpy(m_data,m_cudaData, sizeTex, cudaMemcpyDeviceToHost); + cudaFree(m_cudaData); break; } + // load texture onto the cube + loadTexture(); + m_raytracingNeeded = false; } From b6323fea7ae249fec861799635b44c5e30f39333 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:27:47 +0200 Subject: [PATCH 49/87] add div_up function --- hscgv_uebung_5/raytracer.cu | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index f386e8e..94d6622 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -222,6 +222,15 @@ Raytracer::initCuda() { initPropertiesKernel<<<1,1>>>(d_objList, d_objPropList, g_objectList.size(), d_lightList, d_lightPropList, g_lightList.size()); } +//------------------------------------------------------------------------------------------------- +// Round (a) up to the nearest multiple of (b), then divide by (b) +// + +int div_up(int a, int b) +{ + return (a + b - 1) / b; +} + //! start the rendering routine on the device void Raytracer::renderCuda(float *renderedScene, int xRes, int yRes) From 0f0268bdcea760c9b511852cadc2e029a8e1ff69 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:28:16 +0200 Subject: [PATCH 50/87] rename renderedScene -> cudaData --- hscgv_uebung_5/raytracer.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 94d6622..6679b5a 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -233,7 +233,7 @@ int div_up(int a, int b) //! start the rendering routine on the device void -Raytracer::renderCuda(float *renderedScene, int xRes, int yRes) +Raytracer::renderCuda(float *cudaData, int xRes, int yRes) { // RAY TRACING: dim3 block(32, 16, 1); From f2417a8e1c884e94270981c5294ad68d5ecbb7cf Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:29:03 +0200 Subject: [PATCH 51/87] do the initialization or die tryin --- hscgv_uebung_5/raytracer.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 6679b5a..95996bd 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -239,6 +239,7 @@ Raytracer::renderCuda(float *cudaData, int xRes, int yRes) dim3 block(32, 16, 1); dim3 grid(xRes/ block.x, yRes / block.y, 1); renderKernel<<>>(d_renderedScene, xRes, yRes, + initCuda(); g_scene.view.eyepoint, g_scene.view.up, g_scene.view.lookat, g_scene.view.aspect, g_scene.view.fovy, g_scene.picture.background, m_antialiasing, d_objList, g_objectList.size(), d_lightList, g_lightList.size()); } From be9139464e401fd148c078e08c3d3983cc129cf3 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:32:48 +0200 Subject: [PATCH 52/87] get correct pixel position in grid --- hscgv_uebung_5/raytracer.cu | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 95996bd..736f833 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -134,12 +134,10 @@ renderKernel(float *d_renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d u GeoObject* d_objList, int objListSize, LightObject* d_lightList, int lightListSize) { // find out id of this thread - unsigned int idx = blockIdx.x*blockDim.x + threadIdx.x; - if (idx>=xRes*yRes) + unsigned sx = blockIdx.x * blockDim.x + threadIdx.x; + unsigned sy = blockIdx.y * blockDim.y + threadIdx.y; + if (sx >= xRes || sy >= yRes) return; - // we are handling pixel (x,y) now - int sx = idx % xRes; - int sy = idx / xRes; // setup viewport, its origin is bottom left // setup camera coordsys From f3c05009f18120e6f6aad8f1da45b209846a11a8 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 11:33:46 +0200 Subject: [PATCH 53/87] correct call of kernel (32 are too much) --- hscgv_uebung_5/raytracer.cu | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 736f833..fe23825 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -234,10 +234,11 @@ void Raytracer::renderCuda(float *cudaData, int xRes, int yRes) { // RAY TRACING: - dim3 block(32, 16, 1); - dim3 grid(xRes/ block.x, yRes / block.y, 1); - renderKernel<<>>(d_renderedScene, xRes, yRes, initCuda(); + dim3 block(16, 16); + dim3 grid(div_up(xRes, block.x), div_up(yRes, block.y)); + + renderKernel<<>>(cudaData, xRes, yRes, g_scene.view.eyepoint, g_scene.view.up, g_scene.view.lookat, g_scene.view.aspect, g_scene.view.fovy, g_scene.picture.background, m_antialiasing, d_objList, g_objectList.size(), d_lightList, g_lightList.size()); } From fc63af28847ce186bc715cface8399bacb697e8d Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Mon, 6 Jul 2015 15:53:06 +0200 Subject: [PATCH 54/87] free data later --- hscgv_uebung_5/GLFrame.cpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 8830cea..fa7ade5 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -101,9 +101,6 @@ void GLFrame::loadTexture() // register this texture with CUDA } - // free data since its on the texture - free(m_data); - // cudaGLMapBufferObject((void**)&m_currentDevice, m_currentVBO_CUDA); glBindTexture(GL_TEXTURE_2D, 0); @@ -563,10 +560,10 @@ void GLFrame::renderScene() return; // draw scene int sizeTex = sizeof(float) * 3 * m_width * m_height; - m_data = (float*)malloc(sizeTex); switch(m_renderMode) { case CPU: + m_data = (float*)malloc(sizeTex); m_raytracer->render(m_data, m_width, m_height); break; case GPU: From dd1d3bf4fa01434a9d9c127595c208a97e3b9a64 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 11:01:50 +0200 Subject: [PATCH 55/87] raytracer gets its own structure - remove all host device clutter - setup basic objects as structs - write only-raytracing functions --- hscgv_uebung_5/geoobject.h | 26 +++--- hscgv_uebung_5/geoobject.inl | 20 ++-- hscgv_uebung_5/geoquadric.h | 2 +- hscgv_uebung_5/geoquadric.inl | 10 +- hscgv_uebung_5/ray.inl | 74 +-------------- hscgv_uebung_5/raytracer.cu | 168 +++++++++++++++++++++++++++++++--- hscgv_uebung_5/raytracer.h | 28 +++++- 7 files changed, 211 insertions(+), 117 deletions(-) diff --git a/hscgv_uebung_5/geoobject.h b/hscgv_uebung_5/geoobject.h index 18cfa46..38b7fdb 100644 --- a/hscgv_uebung_5/geoobject.h +++ b/hscgv_uebung_5/geoobject.h @@ -33,11 +33,11 @@ class GeoObjectProperties ~GeoObjectProperties() {} // access methods - Color __host__ __device__ ambient() const; - Vec3d __host__ __device__ reflectance() const; - double __host__ __device__ specular() const; - int __host__ __device__ specularExp() const; - double __host__ __device__ mirror() const; + Color ambient() const; + Vec3d reflectance() const; + double specular() const; + int specularExp() const; + double mirror() const; protected: Color m_ambient; @@ -57,19 +57,19 @@ class GeoObject // CONSTRUCTORS GeoObject() {} virtual ~GeoObject() {} - virtual Vec3d __host__ __device__ getNormal(const Vec3d &v) const = 0; - virtual double __host__ __device__ intersect(const Ray &r) const = 0; + virtual Vec3d getNormal(const Vec3d &v) const = 0; + virtual double intersect(const Ray &r) const = 0; // access methods - virtual Color __host__ __device__ ambient() const; - virtual Vec3d __host__ __device__ reflectance() const; - virtual double __host__ __device__ specular() const; - virtual double __host__ __device__ specularExp() const; + virtual Color ambient() const; + virtual Vec3d reflectance() const; + virtual double specular() const; + virtual double specularExp() const; - virtual double __host__ __device__ mirror() const; + virtual double mirror() const; // config method - virtual void __host__ __device__ setProperties(GeoObjectProperties *p); + virtual void setProperties(GeoObjectProperties *p); protected: GeoObjectProperties *m_properties; diff --git a/hscgv_uebung_5/geoobject.inl b/hscgv_uebung_5/geoobject.inl index fd5ba68..e5b4df2 100644 --- a/hscgv_uebung_5/geoobject.inl +++ b/hscgv_uebung_5/geoobject.inl @@ -8,31 +8,31 @@ /* member access functions for GeoObjectProperties */ -inline Color __host__ __device__ +inline Color GeoObjectProperties::ambient() const { return m_ambient; } -inline Vec3d __host__ __device__ +inline Vec3d GeoObjectProperties::reflectance() const { return m_reflectance; } -inline double __host__ __device__ +inline double GeoObjectProperties::specular() const { return m_specular; } -inline int __host__ __device__ +inline int GeoObjectProperties::specularExp() const { return m_specularExp; } -inline double __host__ __device__ +inline double GeoObjectProperties::mirror() const { return m_mirror; @@ -50,7 +50,7 @@ GeoObject::ambient() const return Vec3d(0.0); } -inline Vec3d __host__ __device__ +inline Vec3d GeoObject::reflectance() const { if (m_properties) @@ -60,7 +60,7 @@ GeoObject::reflectance() const return Vec3d(0.0); } -inline double __host__ __device__ +inline double GeoObject::specular() const { if (m_properties) @@ -70,7 +70,7 @@ GeoObject::specular() const return 0.0; } -inline double __host__ __device__ +inline double GeoObject::specularExp() const { if (m_properties) @@ -80,7 +80,7 @@ GeoObject::specularExp() const return 0.0; } -inline double __host__ __device__ +inline double GeoObject::mirror() const { if (m_properties) @@ -90,7 +90,7 @@ GeoObject::mirror() const return 0.0; } -inline void __host__ __device__ +inline void GeoObject::setProperties(GeoObjectProperties *p) { m_properties = p; diff --git a/hscgv_uebung_5/geoquadric.h b/hscgv_uebung_5/geoquadric.h index 6b594e0..5d1d719 100644 --- a/hscgv_uebung_5/geoquadric.h +++ b/hscgv_uebung_5/geoquadric.h @@ -34,7 +34,7 @@ class GeoQuadric : public GeoObject //! Compute surface normal on the quadric in the point v. virtual Vec3d getNormal(const Vec3d &v) const; //! Compute intersection of ray with quadric. - virtual double __host__ __device__ intersect(const Ray &r) const; + virtual double intersect(const Ray &r) const; private: //! parameters of the equation describing the quadric. diff --git a/hscgv_uebung_5/geoquadric.inl b/hscgv_uebung_5/geoquadric.inl index 835022b..589d175 100644 --- a/hscgv_uebung_5/geoquadric.inl +++ b/hscgv_uebung_5/geoquadric.inl @@ -10,7 +10,7 @@ // Description: // Constructor. - __host__ __device__ GeoQuadric::GeoQuadric() +GeoQuadric::GeoQuadric() : m_a(0), m_b(0), m_c(0) , m_d(0), m_e(0), m_f(0) , m_g(0), m_h(0), m_j(0) @@ -21,7 +21,7 @@ // Description: // explicit parametrization - __host__ __device__ GeoQuadric::GeoQuadric(double na, double nb, double nc, double nd, double ne, +GeoQuadric::GeoQuadric(double na, double nb, double nc, double nd, double ne, double nf, double ng, double nh, double nj, double nk) : m_a(na), m_b(nb), m_c(nc) , m_d(nd), m_e(ne), m_f(nf) @@ -33,14 +33,14 @@ // Description: // Destructor. - __host__ __device__ GeoQuadric::~GeoQuadric() + GeoQuadric::~GeoQuadric() { } // Description: // Return the normal vector at point v of this surface -Vec3d __host__ __device__ +Vec3d GeoQuadric::getNormal(const Vec3d &v) const { Vec3d tmpvec( (v | Vec3d(2*m_a,m_b,m_c)) + m_d, @@ -52,7 +52,7 @@ GeoQuadric::getNormal(const Vec3d &v) const // Description: // Compute intersection point on this surface // return distance to nearest intersection found or -1 -double __host__ __device__ +double GeoQuadric::intersect(const Ray &ray) const { double t = -1.0, acoef, bcoef, ccoef, root, disc; diff --git a/hscgv_uebung_5/ray.inl b/hscgv_uebung_5/ray.inl index a829557..29bc991 100644 --- a/hscgv_uebung_5/ray.inl +++ b/hscgv_uebung_5/ray.inl @@ -64,7 +64,7 @@ inline Ray::~Ray( ) // Description: // Determine color contribution of a lightsource -inline Color __device__ __host__ +inline Color shadedColor(LightObject *light, const Ray &reflectedRay, const Vec3d &normal, GeoObject *obj) { double ldot = light->direction() | normal; @@ -90,7 +90,7 @@ shadedColor(LightObject *light, const Ray &reflectedRay, const Vec3d &normal, Ge // Description: // Determine color of this ray by tracing through the scene -inline const Color __host__ +inline const Color Ray::shade() const { GeoObject *closest = NULL; @@ -162,79 +162,13 @@ Ray::shade() const } } -inline const Color __device__ Ray::shade(Ray *thisRay, Vec3d d_origin, Vec3d d_direction, GeoObject *d_objList, int objListSize, LightObject *d_lightList, int lightListSize, Color background) -{ - Color currentColor(0.0); - for (int i=0; i<5; i++) { - GeoObject *closest = NULL; - double tMin = DBL_MAX; - - // find closest object that intersects - for (int j=0; j ray goes to infinity - if (closest == NULL) { - if (i == 0) { - return background; // background color - } - else { - return Color(0.0); // black - } - } - else { - // reflection - Vec3d intersectionPosition(d_origin + (d_direction * tMin)); - Vec3d normal(closest->getNormal(intersectionPosition)); - Ray reflectedRay(intersectionPosition, - d_direction.getReflectedAt(normal).getNormalized(), - i+1); - - // calculate lighting - for (int j=0; j 0.0) { - something_intersected = true; - break; - } - - } // for all obj - - // is it visible ? - if (! something_intersected) - currentColor += shadedColor(&d_lightList[j], reflectedRay, normal, closest); - - } // for all lights - - // could be right... - currentColor *= closest->mirror(); - } - } - return Color(currentColor); -} - - -inline Vec3d __device__ __host__ +inline Vec3d Ray::origin() const { return Vec3d(m_origin); } -inline Vec3d __device__ __host__ +inline Vec3d Ray::direction() const { return Vec3d(m_direction); diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index fe23825..4821369 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -22,6 +22,149 @@ ReadScene(const char *inputfile) closeSceneFile(); } +// Description: +// Return the normal vector at point v of this surface +// TODO this is an object method, update such that the called on object is included +Vec3d __device__ +cudaGetNormal(const Vec3d &v, const QUADRIC &q) +{ + Vec3d tmpvec( (v | Vec3d(2*q.m_a,q.m_b,q.m_c)) + q.m_d, + (v | Vec3d(q.m_b,2*q.m_e,q.m_f)) + q.m_g, + (v | Vec3d(q.m_c,q.m_f,2*q.m_h)) + q.m_j ); + return tmpvec.getNormalized(); +} +// Description: +// Compute intersection point on this surface +// return distance to nearest intersection found or -1 +// TODO this is an object method, update such that the called on object is included +double __device__ +cudaIntersect(const RAY &ray, QUADRIC &q) +{ + double t = -1.0, acoef, bcoef, ccoef, root, disc; + + acoef = Vec3d( (ray.m_direction | Vec3d(q.m_a,q.m_b,q.m_c)), + q.m_e * ray.m_direction[1] + q.m_f * ray.m_direction[2], + q.m_h * ray.m_direction[2]) + | ray.m_direction; + + bcoef = (Vec3d(q.m_d,q.m_g,q.m_j) | ray.m_direction) + + (ray.m_origin | Vec3d((ray.m_direction | Vec3d(2*q.m_a,q.m_b,q.m_c)), + (ray.m_direction | Vec3d(q.m_b,2*q.m_e,q.m_f)), + (ray.m_direction | Vec3d(q.m_c,q.m_f,2*q.m_h)))); + + ccoef = (ray.m_origin | Vec3d((Vec3d(q.m_a,q.m_b,q.m_c) | ray.m_origin) + q.m_d, + q.m_e * ray.m_origin[1] + q.m_f * ray.m_origin[2] + q.m_g, + q.m_h * ray.m_origin[2] + q.m_j)) + + q.m_k; + + if (acoef != 0.0) { + disc = bcoef * bcoef - 4.0 * acoef * ccoef; + if (disc > -Vec3d::getEpsilon()) { + root = sqrt( disc ); + t = ( -bcoef - root ) / ( acoef + acoef ); + if (t < 0.0) + t = ( -bcoef + root ) / ( acoef + acoef ); + } + } + return (((Vec3d::getEpsilon() * 10.0) < t) ? t : -1.0); +} +// Description: +// Determine color contribution of a lightsource +Color __device__ +cudaShadedColor(LIGHT *light, const RAY &reflectedRay, const Vec3d &normal, QUADRIC *obj) +{ + double ldot = light->m_direction | normal; + Color reflectedColor = Color(0.0); + + // lambertian reflection model + if (ldot > 0.0) + reflectedColor += obj->m_reflectance * (light->m_color * ldot); + + // updated with ambient lightning as in: + // [GENERALISED AMBIENT REFLECTION MODELS FOR LAMBERTIAN AND PHONG SURFACES, Xiaozheng Zhang and Yongsheng Gao] +// reflectedColor += obj->ambient() * g_sceneCuda.ambience; + + // specular part + double spec = reflectedRay.m_direction | light->m_direction; + if (spec > 0.0) { + spec = obj->m_specular * pow(spec, obj->m_specularExp); + reflectedColor += light->m_color * spec; + } + + return Color(reflectedColor); +} + + +Color __device__ +cudaShade(RAY *thisRay, Vec3d d_origin, Vec3d d_direction, QUADRIC *d_objList, int objListSize, LIGHT *d_lightList, int lightListSize, Color background) +{ + Color currentColor(0.0); + for (int i=0; i<5; i++) { + QUADRIC *closest = NULL; + double tMin = DBL_MAX; + + // find closest object that intersects + for (int j=0; j ray goes to infinity + if (closest == NULL) { + if (i == 0) { + return background; // background color + } + else { + return Color(0.0); // black + } + } + else { + // reflection + Vec3d intersectionPosition(d_origin + (d_direction * tMin)); + Vec3d normal(cudaGetNormal(intersectionPosition, *closest)); + RAY reflectedRay; + reflectedRay.m_origin = intersectionPosition; + reflectedRay.m_direction = d_direction.getReflectedAt(normal).getNormalized(); + reflectedRay.m_depth = i+1; + + // calculate lighting + for (int j=0; j 0.0) { + something_intersected = true; + break; + } + + } // for all obj + + // is it visible ? + if (! something_intersected) + currentColor += cudaShadedColor(&d_lightList[j], reflectedRay, normal, closest); + + } // for all lights + + // could be right... + currentColor *= closest->m_mirror; + } + } + return Color(currentColor); +} + Raytracer::Raytracer(): m_isFileLoaded(false), @@ -120,18 +263,9 @@ inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=t } } -void __global__ -initPropertiesKernel(GeoObject* d_objList, GeoObjectProperties* d_objPropList, int d_objListSize, LightObject* d_lightList, LightObjectProperties* d_lightPropList, int d_lightListSize) { - // setup the objects properties one by one, since we needed to copy them by hand into different lists (objPropList and lightPropList) - for (int i=0; i #endif #include "geoobject.h" +#include "geoquadric.h" #include "lightobject.h" #include "ray.h" #include "param.h" @@ -20,6 +21,26 @@ #define M_PI 3.1415927 #endif +struct QUADRIC { + Color m_ambient; + Vec3d m_reflectance; + double m_specular; + int m_specularExp; + double m_mirror; + double m_a,m_b,m_c,m_d,m_e,m_f,m_g,m_h,m_j,m_k; +}; + +struct RAY { + Vec3d m_origin; + Vec3d m_direction; + unsigned int m_depth; +}; + +struct LIGHT { + Color m_color; + Vec3d m_direction; +}; + class Raytracer { public: @@ -33,10 +54,9 @@ class Raytracer bool m_isFileLoaded; //! CUDA Properties - GeoObject* d_objList; - GeoObjectProperties* d_objPropList; - LightObject* d_lightList; - LightObjectProperties* d_lightPropList; + + QUADRIC* d_objList; + LIGHT* d_lightList; private: const char *m_filename; From f5c36fe54d3e8bf345642e8cf863d54250ab5a16 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 11:02:03 +0200 Subject: [PATCH 56/87] make the parameters of a quadric accessible --- hscgv_uebung_5/geoquadric.h | 1 - 1 file changed, 1 deletion(-) diff --git a/hscgv_uebung_5/geoquadric.h b/hscgv_uebung_5/geoquadric.h index 5d1d719..a724597 100644 --- a/hscgv_uebung_5/geoquadric.h +++ b/hscgv_uebung_5/geoquadric.h @@ -36,7 +36,6 @@ class GeoQuadric : public GeoObject //! Compute intersection of ray with quadric. virtual double intersect(const Ray &r) const; - private: //! parameters of the equation describing the quadric. double m_a,m_b,m_c,m_d,m_e,m_f,m_g,m_h,m_j,m_k; }; From de3d317fc03dd42414f10644d876d0016fe676cf Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 11:02:43 +0200 Subject: [PATCH 57/87] less stuff to allocate on GPU --- hscgv_uebung_5/raytracer.cu | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 4821369..52e53c7 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -347,17 +347,15 @@ renderKernel(float *d_renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d u void Raytracer::initCuda() { // get some space for the objects and their properties ( - gpuErrchk (cudaMalloc((void **) &d_objList, sizeof(GeoObject) * g_objectList.size())); - gpuErrchk (cudaMalloc((void **) &d_objPropList, sizeof(GeoObjectProperties) * g_objectList.size())); - gpuErrchk (cudaMalloc((void **) &d_lightList, sizeof(LightObject) * g_lightList.size())); - gpuErrchk (cudaMalloc((void **) &d_lightPropList, sizeof(LightObjectProperties) * g_lightList.size())); - gpuErrchk (cudaMemcpy(d_objList, g_objectList.data(), sizeof(GeoObject) * g_objectList.size(), cudaMemcpyHostToDevice)); gpuErrchk (cudaMemcpy(d_objPropList, g_objectList.data(), sizeof(GeoObjectProperties) * g_objectList.size(), cudaMemcpyHostToDevice)); gpuErrchk (cudaMemcpy(d_lightList, g_objectList.data(), sizeof(LightObject) * g_lightList.size(), cudaMemcpyHostToDevice)); gpuErrchk (cudaMemcpy(d_lightPropList, g_objectList.data(), sizeof(LightObjectProperties) * g_lightList.size(), cudaMemcpyHostToDevice)); initPropertiesKernel<<<1,1>>>(d_objList, d_objPropList, g_objectList.size(), d_lightList, d_lightPropList, g_lightList.size()); + gpuErrchk (cudaMalloc((void **) &d_objList, sizeof(QUADRIC) * g_objectList.size())); + gpuErrchk (cudaMalloc((void **) &d_lightList, sizeof(LIGHT) * g_lightList.size())); + } //------------------------------------------------------------------------------------------------- From 03227cdef2a5fa5818c32e6857329c862848d96b Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 11:07:10 +0200 Subject: [PATCH 58/87] less memcpy needed --- hscgv_uebung_5/raytracer.cu | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 52e53c7..1ad8136 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -347,15 +347,12 @@ renderKernel(float *d_renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d u void Raytracer::initCuda() { // get some space for the objects and their properties ( - gpuErrchk (cudaMemcpy(d_objList, g_objectList.data(), sizeof(GeoObject) * g_objectList.size(), cudaMemcpyHostToDevice)); - gpuErrchk (cudaMemcpy(d_objPropList, g_objectList.data(), sizeof(GeoObjectProperties) * g_objectList.size(), cudaMemcpyHostToDevice)); - gpuErrchk (cudaMemcpy(d_lightList, g_objectList.data(), sizeof(LightObject) * g_lightList.size(), cudaMemcpyHostToDevice)); - gpuErrchk (cudaMemcpy(d_lightPropList, g_objectList.data(), sizeof(LightObjectProperties) * g_lightList.size(), cudaMemcpyHostToDevice)); - initPropertiesKernel<<<1,1>>>(d_objList, d_objPropList, g_objectList.size(), d_lightList, d_lightPropList, g_lightList.size()); gpuErrchk (cudaMalloc((void **) &d_objList, sizeof(QUADRIC) * g_objectList.size())); gpuErrchk (cudaMalloc((void **) &d_lightList, sizeof(LIGHT) * g_lightList.size())); + gpuErrchk (cudaMemcpy(d_objList, quads.data(), sizeof(QUADRIC) * g_objectList.size(), cudaMemcpyHostToDevice)); + gpuErrchk (cudaMemcpy(d_lightList, lights.data(), sizeof(LIGHT) * g_lightList.size(), cudaMemcpyHostToDevice)); } //------------------------------------------------------------------------------------------------- From 09f32cd88beed0b303091511bdb6d38335321e6f Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 11:18:49 +0200 Subject: [PATCH 59/87] new initialization of to be transferred objects --- hscgv_uebung_5/raytracer.cu | 27 ++++++++++++++++++++++++++- 1 file changed, 26 insertions(+), 1 deletion(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 1ad8136..43ce113 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -347,10 +347,35 @@ renderKernel(float *d_renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d u void Raytracer::initCuda() { // get some space for the objects and their properties ( - initPropertiesKernel<<<1,1>>>(d_objList, d_objPropList, g_objectList.size(), d_lightList, d_lightPropList, g_lightList.size()); gpuErrchk (cudaMalloc((void **) &d_objList, sizeof(QUADRIC) * g_objectList.size())); gpuErrchk (cudaMalloc((void **) &d_lightList, sizeof(LIGHT) * g_lightList.size())); + std::vectorquads; + std::vectorlights; + // prepare the light- and geoobjects for transfer to GPU + for (int i=0; im_a = gq.m_a; + tmp->m_b = gq.m_b; + tmp->m_c = gq.m_c; + tmp->m_d = gq.m_d; + tmp->m_e = gq.m_e; + tmp->m_f = gq.m_f; + tmp->m_g = gq.m_g; + tmp->m_h = gq.m_h; + tmp->m_ambient = gq.ambient(); + tmp->m_mirror = gq.mirror(); + tmp->m_reflectance = gq.reflectance(); + tmp->m_specular = gq.specular(); + tmp->m_specularExp = gq.specularExp(); + quads.push_back(tmp); + LIGHT *tmp2; + LightObject lo = g_lightList.at(i); + tmp2->m_color = lo.color(); + tmp2->m_direction = lo.direction(); + lights.push_back(tmp2); + } gpuErrchk (cudaMemcpy(d_objList, quads.data(), sizeof(QUADRIC) * g_objectList.size(), cudaMemcpyHostToDevice)); gpuErrchk (cudaMemcpy(d_lightList, lights.data(), sizeof(LIGHT) * g_lightList.size(), cudaMemcpyHostToDevice)); } From 186decc7ec22549d78d0d73320fd2c4a38dfab52 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 11:38:48 +0200 Subject: [PATCH 60/87] remove host/device and inline code --- .../{geoobject.inl => geoobject.cpp} | 22 +++++++++---------- hscgv_uebung_5/geoobject.h | 1 - .../{geoquadric.inl => geoquadric.cpp} | 0 hscgv_uebung_5/geoquadric.h | 1 - .../{lightobject.inl => lightobject.cpp} | 12 +++++----- hscgv_uebung_5/lightobject.h | 7 +++--- hscgv_uebung_5/{ray.inl => ray.cpp} | 18 +++++++-------- hscgv_uebung_5/ray.h | 19 ++++++++-------- hscgv_uebung_5/src.pro | 8 +++---- 9 files changed, 42 insertions(+), 46 deletions(-) rename hscgv_uebung_5/{geoobject.inl => geoobject.cpp} (91%) rename hscgv_uebung_5/{geoquadric.inl => geoquadric.cpp} (100%) rename hscgv_uebung_5/{lightobject.inl => lightobject.cpp} (81%) rename hscgv_uebung_5/{ray.inl => ray.cpp} (92%) diff --git a/hscgv_uebung_5/geoobject.inl b/hscgv_uebung_5/geoobject.cpp similarity index 91% rename from hscgv_uebung_5/geoobject.inl rename to hscgv_uebung_5/geoobject.cpp index e5b4df2..a0e48f5 100644 --- a/hscgv_uebung_5/geoobject.inl +++ b/hscgv_uebung_5/geoobject.cpp @@ -8,31 +8,31 @@ /* member access functions for GeoObjectProperties */ -inline Color +Color GeoObjectProperties::ambient() const { return m_ambient; } -inline Vec3d +Vec3d GeoObjectProperties::reflectance() const { return m_reflectance; } -inline double +double GeoObjectProperties::specular() const { return m_specular; } -inline int +int GeoObjectProperties::specularExp() const { return m_specularExp; } -inline double +double GeoObjectProperties::mirror() const { return m_mirror; @@ -40,7 +40,7 @@ GeoObjectProperties::mirror() const /* access functions for the properties of a GeoObject */ -inline Color __host__ +Color GeoObject::ambient() const { if (m_properties) @@ -50,7 +50,7 @@ GeoObject::ambient() const return Vec3d(0.0); } -inline Vec3d +Vec3d GeoObject::reflectance() const { if (m_properties) @@ -60,7 +60,7 @@ GeoObject::reflectance() const return Vec3d(0.0); } -inline double +double GeoObject::specular() const { if (m_properties) @@ -70,7 +70,7 @@ GeoObject::specular() const return 0.0; } -inline double +double GeoObject::specularExp() const { if (m_properties) @@ -80,7 +80,7 @@ GeoObject::specularExp() const return 0.0; } -inline double +double GeoObject::mirror() const { if (m_properties) @@ -90,7 +90,7 @@ GeoObject::mirror() const return 0.0; } -inline void +void GeoObject::setProperties(GeoObjectProperties *p) { m_properties = p; diff --git a/hscgv_uebung_5/geoobject.h b/hscgv_uebung_5/geoobject.h index 38b7fdb..fe1f07c 100644 --- a/hscgv_uebung_5/geoobject.h +++ b/hscgv_uebung_5/geoobject.h @@ -74,6 +74,5 @@ class GeoObject protected: GeoObjectProperties *m_properties; }; -#include "geoobject.inl" #endif /* GEOOBJECT_HH */ diff --git a/hscgv_uebung_5/geoquadric.inl b/hscgv_uebung_5/geoquadric.cpp similarity index 100% rename from hscgv_uebung_5/geoquadric.inl rename to hscgv_uebung_5/geoquadric.cpp diff --git a/hscgv_uebung_5/geoquadric.h b/hscgv_uebung_5/geoquadric.h index a724597..9942ad2 100644 --- a/hscgv_uebung_5/geoquadric.h +++ b/hscgv_uebung_5/geoquadric.h @@ -39,6 +39,5 @@ class GeoQuadric : public GeoObject //! parameters of the equation describing the quadric. double m_a,m_b,m_c,m_d,m_e,m_f,m_g,m_h,m_j,m_k; }; -#include "geoquadric.inl" #endif /* GEOQUADRIC_HH */ diff --git a/hscgv_uebung_5/lightobject.inl b/hscgv_uebung_5/lightobject.cpp similarity index 81% rename from hscgv_uebung_5/lightobject.inl rename to hscgv_uebung_5/lightobject.cpp index d168807..b8fffef 100644 --- a/hscgv_uebung_5/lightobject.inl +++ b/hscgv_uebung_5/lightobject.cpp @@ -8,14 +8,14 @@ // Description: // Constructor. -inline LightObject::LightObject() +LightObject::LightObject() : m_properties(NULL) { } // Description: // explicit parametrization -inline LightObject::LightObject(const Vec3d& dir, const Color& col) +LightObject::LightObject(const Vec3d& dir, const Color& col) { m_properties = new LightObjectProperties(); m_properties->direction = dir.getNormalized(); @@ -24,12 +24,12 @@ inline LightObject::LightObject(const Vec3d& dir, const Color& col) // Description: // Destructor. -inline LightObject::~LightObject() +LightObject::~LightObject() { delete m_properties; } -inline __device__ __host__ Color +Color LightObject::color() const { if (m_properties) @@ -39,7 +39,7 @@ LightObject::color() const return Color(0.0); } -inline __device__ __host__ Vec3d +Vec3d LightObject::direction() const { if (m_properties) @@ -49,7 +49,7 @@ LightObject::direction() const return Vec3d(0.0); } -inline void __host__ __device__ +void LightObject::setProperties(LightObjectProperties *p) { m_properties = p; diff --git a/hscgv_uebung_5/lightobject.h b/hscgv_uebung_5/lightobject.h index ba959fa..0490de2 100644 --- a/hscgv_uebung_5/lightobject.h +++ b/hscgv_uebung_5/lightobject.h @@ -29,14 +29,13 @@ class LightObject virtual ~LightObject(); // access methods - Color __device__ __host__ color() const; - Vec3d __device__ __host__ direction() const; - void __host__ __device__ setProperties(LightObjectProperties *p); + Color color() const; + Vec3d direction() const; + void setProperties(LightObjectProperties *p); protected: LightObjectProperties *m_properties; }; -#include "lightobject.inl" #endif /* LIGHTOBJECT_HH */ diff --git a/hscgv_uebung_5/ray.inl b/hscgv_uebung_5/ray.cpp similarity index 92% rename from hscgv_uebung_5/ray.inl rename to hscgv_uebung_5/ray.cpp index 29bc991..27fe180 100644 --- a/hscgv_uebung_5/ray.inl +++ b/hscgv_uebung_5/ray.cpp @@ -13,7 +13,7 @@ // Description: // Constructor. -inline Ray::Ray( ) +Ray::Ray( ) : m_origin() , m_direction() , m_depth(0) @@ -24,7 +24,7 @@ inline Ray::Ray( ) // Description: // Copy-Constructor. -inline Ray::Ray( const Ray &r ) +Ray::Ray( const Ray &r ) { // copy it ! initialization is not enough ! m_origin = r.m_origin; @@ -36,7 +36,7 @@ inline Ray::Ray( const Ray &r ) // Description: // Constructor with explicit parameters -inline Ray::Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll) +Ray::Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll) { // copy it ! initialization is not enough ! m_origin = o; @@ -48,7 +48,7 @@ inline Ray::Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vectordirection() | normal; @@ -90,7 +90,7 @@ shadedColor(LightObject *light, const Ray &reflectedRay, const Vec3d &normal, Ge // Description: // Determine color of this ray by tracing through the scene -inline const Color +const Color Ray::shade() const { GeoObject *closest = NULL; @@ -162,13 +162,13 @@ Ray::shade() const } } -inline Vec3d +Vec3d Ray::origin() const { return Vec3d(m_origin); } -inline Vec3d +Vec3d Ray::direction() const { return Vec3d(m_direction); diff --git a/hscgv_uebung_5/ray.h b/hscgv_uebung_5/ray.h index bd9a5fc..e5c5b43 100644 --- a/hscgv_uebung_5/ray.h +++ b/hscgv_uebung_5/ray.h @@ -23,18 +23,18 @@ class Ray { public: // CONSTRUCTORS - __host__ __device__ Ray(); - __host__ __device__ Ray(const Ray &r); - __host__ Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll); - __device__ Ray(const Vec3d o, const Vec3d d, unsigned int i); - __host__ __device__ ~Ray(); + Ray(); + Ray(const Ray &r); + Ray(const Vec3d &o, const Vec3d &d, unsigned int i, std::vector &ol, std::vector &ll); + Ray(const Vec3d o, const Vec3d d, unsigned int i); + ~Ray(); - const Color __host__ shade() const; - const Color __device__ shade(Ray *thisRay, Vec3d d_origin, Vec3d d_direction, GeoObject *d_objList, int objListSize, LightObject *d_lightList, int lightListSize, Color background) ; + const Color shade() const; + const Color shade(Ray *thisRay, Vec3d d_origin, Vec3d d_direction, GeoObject *d_objList, int objListSize, LightObject *d_lightList, int lightListSize, Color background) ; // access methods - Vec3d __host__ __device__ origin() const; - Vec3d __host__ __device__ direction() const; + Vec3d origin() const; + Vec3d direction() const; private: Vec3d m_origin; @@ -44,6 +44,5 @@ class Ray std::vector *m_objList; std::vector *m_lightList; }; -#include "ray.inl" #endif /* RAY_HH */ diff --git a/hscgv_uebung_5/src.pro b/hscgv_uebung_5/src.pro index 685996e..3912e84 100644 --- a/hscgv_uebung_5/src.pro +++ b/hscgv_uebung_5/src.pro @@ -8,10 +8,10 @@ HEADERS= geoobject.h geoquadric.h lightobject.h ray.h types.h vector.h pa SOURCES= main.cpp \ ApplicationWindow.cpp \ GLFrame.cpp \ - ray.inl \ - geoobject.inl \ - geoquadric.inl \ - lightobject.inl + geoquadric.cpp \ + lightobject.cpp \ + ray.cpp \ + geoobject.cpp win32:SOURCES*= xgetopt.cpp YACCSOURCES= input.y CUDA_SOURCES += \ From 8bd1fa30fa0306c734770d99397670a122bbb739 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 11:39:17 +0200 Subject: [PATCH 61/87] quickwins --- hscgv_uebung_5/raytracer.cu | 36 ++++++++++++++++++------------------ 1 file changed, 18 insertions(+), 18 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 43ce113..1bdac05 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -353,27 +353,27 @@ Raytracer::initCuda() { std::vectorquads; std::vectorlights; // prepare the light- and geoobjects for transfer to GPU - for (int i=0; im_a = gq.m_a; - tmp->m_b = gq.m_b; - tmp->m_c = gq.m_c; - tmp->m_d = gq.m_d; - tmp->m_e = gq.m_e; - tmp->m_f = gq.m_f; - tmp->m_g = gq.m_g; - tmp->m_h = gq.m_h; - tmp->m_ambient = gq.ambient(); - tmp->m_mirror = gq.mirror(); - tmp->m_reflectance = gq.reflectance(); - tmp->m_specular = gq.specular(); - tmp->m_specularExp = gq.specularExp(); + GeoQuadric *gq = (GeoQuadric *)g_objectList.at(i); + tmp->m_a = gq->m_a; + tmp->m_b = gq->m_b; + tmp->m_c = gq->m_c; + tmp->m_d = gq->m_d; + tmp->m_e = gq->m_e; + tmp->m_f = gq->m_f; + tmp->m_g = gq->m_g; + tmp->m_h = gq->m_h; + tmp->m_ambient = gq->ambient(); + tmp->m_mirror = gq->mirror(); + tmp->m_reflectance = gq->reflectance(); + tmp->m_specular = gq->specular(); + tmp->m_specularExp = gq->specularExp(); quads.push_back(tmp); LIGHT *tmp2; - LightObject lo = g_lightList.at(i); - tmp2->m_color = lo.color(); - tmp2->m_direction = lo.direction(); + LightObject *lo = g_lightList.at(i); + tmp2->m_color = lo->color(); + tmp2->m_direction = lo->direction(); lights.push_back(tmp2); } gpuErrchk (cudaMemcpy(d_objList, quads.data(), sizeof(QUADRIC) * g_objectList.size(), cudaMemcpyHostToDevice)); From abeb4e10edb30e42f0c730972c2b2ffc01be72bf Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 13:56:11 +0200 Subject: [PATCH 62/87] we don't need pointer --- hscgv_uebung_5/raytracer.cu | 40 +++++++++++++++++++------------------ 1 file changed, 21 insertions(+), 19 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 1bdac05..24c6cab 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -350,30 +350,32 @@ Raytracer::initCuda() { gpuErrchk (cudaMalloc((void **) &d_objList, sizeof(QUADRIC) * g_objectList.size())); gpuErrchk (cudaMalloc((void **) &d_lightList, sizeof(LIGHT) * g_lightList.size())); - std::vectorquads; - std::vectorlights; + std::vector quads; + std::vector lights; // prepare the light- and geoobjects for transfer to GPU for (unsigned int i=0; im_a = gq->m_a; - tmp->m_b = gq->m_b; - tmp->m_c = gq->m_c; - tmp->m_d = gq->m_d; - tmp->m_e = gq->m_e; - tmp->m_f = gq->m_f; - tmp->m_g = gq->m_g; - tmp->m_h = gq->m_h; - tmp->m_ambient = gq->ambient(); - tmp->m_mirror = gq->mirror(); - tmp->m_reflectance = gq->reflectance(); - tmp->m_specular = gq->specular(); - tmp->m_specularExp = gq->specularExp(); + tmp.m_a = gq->m_a; + tmp.m_b = gq->m_b; + tmp.m_c = gq->m_c; + tmp.m_d = gq->m_d; + tmp.m_e = gq->m_e; + tmp.m_f = gq->m_f; + tmp.m_g = gq->m_g; + tmp.m_h = gq->m_h; + tmp.m_ambient = gq->ambient(); + tmp.m_mirror = gq->mirror(); + tmp.m_reflectance = gq->reflectance(); + tmp.m_specular = gq->specular(); + tmp.m_specularExp = gq->specularExp(); quads.push_back(tmp); - LIGHT *tmp2; + } + for (unsigned int i=0; im_color = lo->color(); - tmp2->m_direction = lo->direction(); + tmp2.m_color = lo->color(); + tmp2.m_direction = lo->direction(); lights.push_back(tmp2); } gpuErrchk (cudaMemcpy(d_objList, quads.data(), sizeof(QUADRIC) * g_objectList.size(), cudaMemcpyHostToDevice)); From 996a08644b7eb0e843afab72b9d092eb03af1003 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 13:56:29 +0200 Subject: [PATCH 63/87] add DBL_MAX --- hscgv_uebung_5/raytracer.h | 1 + 1 file changed, 1 insertion(+) diff --git a/hscgv_uebung_5/raytracer.h b/hscgv_uebung_5/raytracer.h index 7ff9120..3908a63 100644 --- a/hscgv_uebung_5/raytracer.h +++ b/hscgv_uebung_5/raytracer.h @@ -16,6 +16,7 @@ #include "ray.h" #include "param.h" #include "types.h" +#include #ifndef M_PI #define M_PI 3.1415927 From 37e68d09efac6fc04aca0c3bf632a2a9a1bddd97 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 15:50:30 +0200 Subject: [PATCH 64/87] initCuda should be called only once a time a model is loaded --- hscgv_uebung_5/raytracer.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 24c6cab..f3a4130 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -396,7 +396,6 @@ void Raytracer::renderCuda(float *cudaData, int xRes, int yRes) { // RAY TRACING: - initCuda(); dim3 block(16, 16); dim3 grid(div_up(xRes, block.x), div_up(yRes, block.y)); From 20283e4500f450be1925db650b0b04881061f2df Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 15:52:17 +0200 Subject: [PATCH 65/87] FIXBUG and rename --- hscgv_uebung_5/GLFrame.cpp | 26 ++++++++++++++++++-------- 1 file changed, 18 insertions(+), 8 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index fa7ade5..2055e25 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -250,6 +250,20 @@ static void matMult(GLfloat *out, const GLfloat *a, const GLfloat *b) memcpy(out, n, sizeof(n)); } +static Vec3d operator *(const GLfloat *a, const Vec3d &b) +{ + // handle case of out==a or out==b + Vec3d n; + + for(int i=0; i<3; i++) { + n[i] = 0.0; + for(int j=0; j<3; j++) { + n[i] += a[i*4+j]*b[j]; + } + } + + return n; +} static void matRot(GLfloat *out, const GLfloat *a, const GLfloat rad) { @@ -395,16 +409,12 @@ void GLFrame::adjustCam(bool leftButton, bool middleButton, bool rightButton, if(d > 0.0001) { float a[4] = {dy/d, -dx/d, 0.0, 0.0}; // rotation axis - GLfloat inc[16]; - matRot(inc, a, d*5); + GLfloat rot[16]; + matRot(rot, a, d*5); // multiply transformation matrix with eyepoint to get new perspective - Vec3d eyeg = Vec3d(g_scene.view.eyepoint); - GLfloat eye[4] = {eyeg[0],eyeg[1],eyeg[2],.0}; - matMult(eye, inc, eye); - for (int i=0;i<3;i++) - g_scene.view.eyepoint[i] = eye[i]; - + Vec3d eyeg = g_scene.view.eyepoint; + g_scene.view.eyepoint = rot * eyeg; } } else if(middleButton) From bc9047e608e4209414091a92074eda07f366bae2 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Tue, 7 Jul 2015 15:52:46 +0200 Subject: [PATCH 66/87] we want to raytrace a lot of times, since its fast --- hscgv_uebung_5/GLFrame.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 2055e25..2481c19 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -586,7 +586,7 @@ void GLFrame::renderScene() } // load texture onto the cube loadTexture(); - m_raytracingNeeded = false; + m_raytracingNeeded = true; } From 5105c7eba0345d5ef043b9f615779842c6877351 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 14:12:31 +0200 Subject: [PATCH 67/87] FIXBUG we need to allocate some place anyway --- hscgv_uebung_5/GLFrame.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 2481c19..3ee2324 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -570,10 +570,10 @@ void GLFrame::renderScene() return; // draw scene int sizeTex = sizeof(float) * 3 * m_width * m_height; + m_data = (float*)malloc(sizeTex); switch(m_renderMode) { case CPU: - m_data = (float*)malloc(sizeTex); m_raytracer->render(m_data, m_width, m_height); break; case GPU: From 1923c11b6627290d6012acc52aece7394b412186 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 14:12:57 +0200 Subject: [PATCH 68/87] remove useful error messages --- hscgv_uebung_5/GLFrame.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 3ee2324..74b6311 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -579,7 +579,6 @@ void GLFrame::renderScene() case GPU: cudaMalloc(&m_cudaData,sizeTex); m_raytracer->renderCuda(m_cudaData, m_width, m_height); - cudaGetLastError(); cudaMemcpy(m_data,m_cudaData, sizeTex, cudaMemcpyDeviceToHost); cudaFree(m_cudaData); break; From 0989ae3e74f15fc1675bc08781e5f2e65523780c Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 14:13:23 +0200 Subject: [PATCH 69/87] rename input --- hscgv_uebung_5/raytracer.cu | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index f3a4130..4d9c13a 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -264,7 +264,7 @@ inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=t } void __global__ -renderKernel(float *d_renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, double fovy, Color backgroundCol, bool antialiasing, +renderKernel(float *d_cudaData, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Vec3d lookat, double aspect, double fovy, Color backgroundCol, bool antialiasing, QUADRIC* d_objList, int objListSize, LIGHT* d_lightList, int lightListSize) { // find out id of this thread @@ -339,6 +339,9 @@ renderKernel(float *d_renderedScene, int xRes, int yRes, Vec3d eyepoint, Vec3d u d_renderedScene[index + 0] = col[0]; d_renderedScene[index + 1] = col[1]; d_renderedScene[index + 2] = col[2]; + d_cudaData[index + 0] = col[0]; + d_cudaData[index + 1] = col[1]; + d_cudaData[index + 2] = col[2]; } From a778e9468d1b4554a5ba6565ef7d9e162e81dbc9 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 14:13:37 +0200 Subject: [PATCH 70/87] FIXBUG I don't use shaded colourpoint --- hscgv_uebung_5/raytracer.cu | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 4d9c13a..4ab0885 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -303,8 +303,7 @@ renderKernel(float *d_cudaData, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Ve theRay.m_depth = 0; // compute the color - Color col; - cudaShade(&theRay, eyepoint, point, d_objList, objListSize, d_lightList, lightListSize, backgroundCol); + Color col= cudaShade(&theRay, eyepoint, point, d_objList, objListSize, d_lightList, lightListSize, backgroundCol); // in case we are using antialiasing, calculate the color of this pixel by averaging if (antialiasing) { From 7404b87e95de4b660b68d46fb66e8d8722627c3f Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 14:14:09 +0200 Subject: [PATCH 71/87] add useful comment --- hscgv_uebung_5/raytracer.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 4ab0885..91a6ed5 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -380,6 +380,7 @@ Raytracer::initCuda() { tmp2.m_direction = lo->direction(); lights.push_back(tmp2); } + // copy objects and lights to GPU gpuErrchk (cudaMemcpy(d_objList, quads.data(), sizeof(QUADRIC) * g_objectList.size(), cudaMemcpyHostToDevice)); gpuErrchk (cudaMemcpy(d_lightList, lights.data(), sizeof(LIGHT) * g_lightList.size(), cudaMemcpyHostToDevice)); } From b66f47801c04039b4303ba7673026f5c625c461d Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 14:15:53 +0200 Subject: [PATCH 72/87] visualize direction vector --- hscgv_uebung_5/raytracer.cu | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 91a6ed5..8720396 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -302,6 +302,12 @@ renderKernel(float *d_cudaData, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Ve theRay.m_direction = dir.getNormalized(); theRay.m_depth = 0; + int index = 3*((sy-1) * xRes + sx); + d_cudaData[index + 0] = theRay.m_direction[0]; + d_cudaData[index + 1] = theRay.m_direction[1]; + d_cudaData[index + 2] = theRay.m_direction[2]; + return; + // compute the color Color col= cudaShade(&theRay, eyepoint, point, d_objList, objListSize, d_lightList, lightListSize, backgroundCol); @@ -334,10 +340,7 @@ renderKernel(float *d_cudaData, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Ve } } - int index = 3*((sy-1) * xRes + sx); - d_renderedScene[index + 0] = col[0]; - d_renderedScene[index + 1] = col[1]; - d_renderedScene[index + 2] = col[2]; +// int index = 3*((sy-1) * xRes + sx); d_cudaData[index + 0] = col[0]; d_cudaData[index + 1] = col[1]; d_cudaData[index + 2] = col[2]; From d6ef7684f0d7db415ea1b4dcecfcd7c16892cfdb Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 17:12:00 +0200 Subject: [PATCH 73/87] FIX slots and menu actions --- hscgv_uebung_5/ApplicationWindow.cpp | 24 ++++-------------------- hscgv_uebung_5/ApplicationWindow.h | 8 +------- hscgv_uebung_5/ApplicationWindow.ui | 10 +++------- hscgv_uebung_5/GLFrame.cpp | 11 ++--------- hscgv_uebung_5/GLFrame.h | 2 +- 5 files changed, 11 insertions(+), 44 deletions(-) diff --git a/hscgv_uebung_5/ApplicationWindow.cpp b/hscgv_uebung_5/ApplicationWindow.cpp index bde899f..74c7e72 100644 --- a/hscgv_uebung_5/ApplicationWindow.cpp +++ b/hscgv_uebung_5/ApplicationWindow.cpp @@ -29,10 +29,9 @@ ApplicationWindow::ApplicationWindow() m_frame = new GLFrame(this); setCentralWidget(m_frame); - QActionGroup *renderModeActions = new QActionGroup(this); - renderModeActions->addAction(ui.actionCPU); +/* QActionGroup *renderModeActions = new QActionGroup(this); renderModeActions->addAction(ui.actionGPU); - renderModeActions->addAction(ui.actionAntialiasing); + renderModeActions->addAction(ui.actionAntialiasing);*/ // connect actions connect(ui.actionQuit, SIGNAL(triggered()), this, SLOT(fileQuit())); @@ -43,8 +42,7 @@ ApplicationWindow::ApplicationWindow() connect(ui.actionAnimate, SIGNAL(toggled(bool)), this, SLOT(animate(bool))); connect(ui.actionAntialiasing, SIGNAL(toggled(bool)), SLOT(antialiasing(bool))); - connect(ui.actionCPU, SIGNAL(toggled(bool)), SLOT(renderModeChanged())); - connect(ui.actionGPU, SIGNAL(toggled(bool)), SLOT(renderModeChanged())); + connect(ui.actionGPU, SIGNAL(toggled(bool)), m_frame, SLOT(setRenderMode(bool))); // ----- SIGNALS ----- @@ -57,7 +55,7 @@ ApplicationWindow::ApplicationWindow() connect(ui.actionLoadFile, SIGNAL(triggered()), this, SLOT(loadFile())); connect(this, SIGNAL(openFile(const QString&)), m_frame, SLOT(loadScene(const QString&))); - connect(this, SIGNAL(renderMode(int)), m_frame, SLOT(setRenderMode(int))); + connect(this, SIGNAL(renderMode(bool)), m_frame, SLOT(setRenderMode(bool))); // timer for updating fps display QTimer *timer = new QTimer(this); @@ -85,8 +83,6 @@ void ApplicationWindow::initState() const ui.actionAnimate->setChecked(true); ui.actionAnimate->trigger(); - - ui.actionCPU->trigger(); } @@ -161,16 +157,4 @@ void ApplicationWindow::antialiasing(bool on) m_frame->m_antialiasing = on; } -int ApplicationWindow::getRenderMode() const -{ - if(ui.actionCPU->isChecked()) - return GLFrame::CPU; - else //if(ui.actionGPU->isChecked()) - return GLFrame::GPU; -} - -void ApplicationWindow::renderModeChanged() const -{ - emit renderMode(getRenderMode()); -} diff --git a/hscgv_uebung_5/ApplicationWindow.h b/hscgv_uebung_5/ApplicationWindow.h index b21c6e5..8b3245d 100644 --- a/hscgv_uebung_5/ApplicationWindow.h +++ b/hscgv_uebung_5/ApplicationWindow.h @@ -60,7 +60,7 @@ class ApplicationWindow : public QMainWindow void openFile(const QString &filename) const; //! render mode was changed - void renderMode(int mode) const; + void renderMode(bool mode) const; protected: //! initialize application state @@ -99,12 +99,6 @@ class ApplicationWindow : public QMainWindow //! we want to have some nice output void antialiasing(bool on); - //! deduce render mode from menu state - int getRenderMode() const; - - //! make sure that render mode signal is triggerd - void renderModeChanged() const; - private: //! store status bar message diff --git a/hscgv_uebung_5/ApplicationWindow.ui b/hscgv_uebung_5/ApplicationWindow.ui index 038d964..9c803ff 100644 --- a/hscgv_uebung_5/ApplicationWindow.ui +++ b/hscgv_uebung_5/ApplicationWindow.ui @@ -42,9 +42,7 @@ Rendering - - @@ -113,12 +111,10 @@ QAction::TextHeuristicRole - - - CPU - - + + true + GPU diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 74b6311..8fa77b8 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -440,16 +440,9 @@ void GLFrame::adjustCam(bool leftButton, bool middleButton, bool rightButton, } // notify when the drawMode should be changed -void GLFrame::setRenderMode(int mode) +void GLFrame::setRenderMode(bool mode) { - m_renderMode = (RenderMode)mode; - - if(mode == CPU) { - // unset VBO - } - else if (mode == GPU) { - } - + m_renderMode = mode ? GPU : CPU;//(RenderMode)mode; m_raytracingNeeded = true; updateGL(); } diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index 8eafa17..8e603af 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -84,7 +84,7 @@ class GLFrame : public QGLWidget public slots: //! change render mode (wireframe, ...) - void setRenderMode(int mode); + void setRenderMode(bool mode); //! set visibility of coordinate axes void setAxesVisibility(bool on); //! set antialising From 4a54a7b60670b3fd1543138d89e70a6d9f6b0de2 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 17:13:17 +0200 Subject: [PATCH 74/87] FIX shading with accumulated mirror factor --- hscgv_uebung_5/raytracer.cu | 9 ++++++--- hscgv_uebung_5/vector.h | 3 +++ 2 files changed, 9 insertions(+), 3 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 8720396..53c3af3 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -99,6 +99,8 @@ Color __device__ cudaShade(RAY *thisRay, Vec3d d_origin, Vec3d d_direction, QUADRIC *d_objList, int objListSize, LIGHT *d_lightList, int lightListSize, Color background) { Color currentColor(0.0); + Color::value_type currentMirror(1.0); + // iterate at most 5 times or until the ray hits the background for (int i=0; i<5; i++) { QUADRIC *closest = NULL; double tMin = DBL_MAX; @@ -153,13 +155,14 @@ cudaShade(RAY *thisRay, Vec3d d_origin, Vec3d d_direction, QUADRIC *d_objList, i } // for all obj // is it visible ? - if (! something_intersected) - currentColor += cudaShadedColor(&d_lightList[j], reflectedRay, normal, closest); + if (!something_intersected) + currentColor += cudaShadedColor(&d_lightList[j], reflectedRay, normal, closest)*currentMirror; + } // for all lights // could be right... - currentColor *= closest->m_mirror; + currentMirror *= closest->m_mirror; } } return Color(currentColor); diff --git a/hscgv_uebung_5/vector.h b/hscgv_uebung_5/vector.h index 4e0be9e..b06ac77 100644 --- a/hscgv_uebung_5/vector.h +++ b/hscgv_uebung_5/vector.h @@ -41,6 +41,9 @@ operator>>( std::istream& is, Vector& i ); template class Vector { + public: + typedef Scalar value_type; + public: //! Constructor __host__ __device__ Vector(void ); From 399265e5bfafdc7661dfe4efb2868732c8533d09 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 17:14:04 +0200 Subject: [PATCH 75/87] FIX call to cudaShade does need only a ray --- hscgv_uebung_5/raytracer.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 53c3af3..58ec405 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -96,7 +96,7 @@ cudaShadedColor(LIGHT *light, const RAY &reflectedRay, const Vec3d &normal, QUAD Color __device__ -cudaShade(RAY *thisRay, Vec3d d_origin, Vec3d d_direction, QUADRIC *d_objList, int objListSize, LIGHT *d_lightList, int lightListSize, Color background) +cudaShade(RAY *thisRay, QUADRIC *d_objList, int objListSize, LIGHT *d_lightList, int lightListSize, Color background) { Color currentColor(0.0); Color::value_type currentMirror(1.0); @@ -312,7 +312,7 @@ renderKernel(float *d_cudaData, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Ve return; // compute the color - Color col= cudaShade(&theRay, eyepoint, point, d_objList, objListSize, d_lightList, lightListSize, backgroundCol); + Color col = cudaShade(&theRay, d_objList, objListSize, d_lightList, lightListSize, backgroundCol); // in case we are using antialiasing, calculate the color of this pixel by averaging if (antialiasing) { @@ -337,7 +337,7 @@ renderKernel(float *d_cudaData, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Ve theRayA.m_origin = eyepoint; theRayA.m_direction = superSampleDir.getNormalized(); theRayA.m_depth = 0; - col += cudaShade(&theRay,eyepoint,superSamplePoint, d_objList, objListSize, d_lightList, lightListSize, backgroundCol)*0.2; + col += cudaShade(&theRay, d_objList, objListSize, d_lightList, lightListSize, backgroundCol)*0.2; //color, recursive_ray_trace(eye, ray, 0)); } } From b0af31380c9533f578dcddb5f4fb6f824f84250b Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 17:14:48 +0200 Subject: [PATCH 76/87] FIXBUG remove this code since we don't use recursion --- hscgv_uebung_5/raytracer.cu | 3 --- 1 file changed, 3 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 58ec405..55bc04e 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -120,9 +120,6 @@ cudaShade(RAY *thisRay, QUADRIC *d_objList, int objListSize, LIGHT *d_lightList, if (i == 0) { return background; // background color } - else { - return Color(0.0); // black - } } else { // reflection From 642d68f99ccbf10087b45478f94dbcd26a82450f Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 17:15:22 +0200 Subject: [PATCH 77/87] FIXBUG used wrong origin and direction --- hscgv_uebung_5/raytracer.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 55bc04e..4f9cbb1 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -123,12 +123,12 @@ cudaShade(RAY *thisRay, QUADRIC *d_objList, int objListSize, LIGHT *d_lightList, } else { // reflection - Vec3d intersectionPosition(d_origin + (d_direction * tMin)); + Vec3d intersectionPosition(thisRay->m_origin + (thisRay->m_direction * tMin)); Vec3d normal(cudaGetNormal(intersectionPosition, *closest)); RAY reflectedRay; reflectedRay.m_origin = intersectionPosition; - reflectedRay.m_direction = d_direction.getReflectedAt(normal).getNormalized(); reflectedRay.m_depth = i+1; + reflectedRay.m_direction = thisRay->m_direction.getReflectedAt(normal).getNormalized(); // calculate lighting for (int j=0; j Date: Wed, 8 Jul 2015 17:15:41 +0200 Subject: [PATCH 78/87] remove unneeded code --- hscgv_uebung_5/raytracer.cu | 5 ----- 1 file changed, 5 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 4f9cbb1..923cfcf 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -80,9 +80,6 @@ cudaShadedColor(LIGHT *light, const RAY &reflectedRay, const Vec3d &normal, QUAD if (ldot > 0.0) reflectedColor += obj->m_reflectance * (light->m_color * ldot); - // updated with ambient lightning as in: - // [GENERALISED AMBIENT REFLECTION MODELS FOR LAMBERTIAN AND PHONG SURFACES, Xiaozheng Zhang and Yongsheng Gao] -// reflectedColor += obj->ambient() * g_sceneCuda.ambience; // specular part double spec = reflectedRay.m_direction | light->m_direction; @@ -127,7 +124,6 @@ cudaShade(RAY *thisRay, QUADRIC *d_objList, int objListSize, LIGHT *d_lightList, Vec3d normal(cudaGetNormal(intersectionPosition, *closest)); RAY reflectedRay; reflectedRay.m_origin = intersectionPosition; - reflectedRay.m_depth = i+1; reflectedRay.m_direction = thisRay->m_direction.getReflectedAt(normal).getNormalized(); // calculate lighting @@ -137,7 +133,6 @@ cudaShade(RAY *thisRay, QUADRIC *d_objList, int objListSize, LIGHT *d_lightList, RAY rayoflight; rayoflight.m_origin = intersectionPosition; rayoflight.m_direction = d_lightList[j].m_direction; - rayoflight.m_depth = 0; bool something_intersected = false; // where are the objects ? From 5924911f8935cbfab89143f86af5f4d176a25090 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 17:17:52 +0200 Subject: [PATCH 79/87] FIXBUG we need to update thisRay --- hscgv_uebung_5/raytracer.cu | 3 +++ 1 file changed, 3 insertions(+) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 923cfcf..451c3b5 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -155,6 +155,9 @@ cudaShade(RAY *thisRay, QUADRIC *d_objList, int objListSize, LIGHT *d_lightList, // could be right... currentMirror *= closest->m_mirror; + // update thisRay to new point + *thisRay = reflectedRay; + } } return Color(currentColor); From e2607a609c02922d723ad99ef198533cc8563d3d Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 17:18:12 +0200 Subject: [PATCH 80/87] add comment --- hscgv_uebung_5/raytracer.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 451c3b5..8250741 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -92,6 +92,7 @@ cudaShadedColor(LIGHT *light, const RAY &reflectedRay, const Vec3d &normal, QUAD } +// Determine color of this ray by tracing through the scene Color __device__ cudaShade(RAY *thisRay, QUADRIC *d_objList, int objListSize, LIGHT *d_lightList, int lightListSize, Color background) { From c338674c1c7149476bae3d53ef819ca1b09cc7e3 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 17:19:01 +0200 Subject: [PATCH 81/87] FIXBUG forgot to upload all quadric attributes --- hscgv_uebung_5/raytracer.cu | 2 ++ 1 file changed, 2 insertions(+) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 8250741..55132a2 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -368,6 +368,8 @@ Raytracer::initCuda() { tmp.m_f = gq->m_f; tmp.m_g = gq->m_g; tmp.m_h = gq->m_h; + tmp.m_j = gq->m_j; + tmp.m_k = gq->m_k; tmp.m_ambient = gq->ambient(); tmp.m_mirror = gq->mirror(); tmp.m_reflectance = gq->reflectance(); From fd9af2704e76e0f8c3e9e3c2a81b84eb31841523 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 17:19:12 +0200 Subject: [PATCH 82/87] undo debug code --- hscgv_uebung_5/raytracer.cu | 9 ++------- 1 file changed, 2 insertions(+), 7 deletions(-) diff --git a/hscgv_uebung_5/raytracer.cu b/hscgv_uebung_5/raytracer.cu index 55132a2..77388bd 100644 --- a/hscgv_uebung_5/raytracer.cu +++ b/hscgv_uebung_5/raytracer.cu @@ -301,12 +301,6 @@ renderKernel(float *d_cudaData, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Ve theRay.m_direction = dir.getNormalized(); theRay.m_depth = 0; - int index = 3*((sy-1) * xRes + sx); - d_cudaData[index + 0] = theRay.m_direction[0]; - d_cudaData[index + 1] = theRay.m_direction[1]; - d_cudaData[index + 2] = theRay.m_direction[2]; - return; - // compute the color Color col = cudaShade(&theRay, d_objList, objListSize, d_lightList, lightListSize, backgroundCol); @@ -339,10 +333,11 @@ renderKernel(float *d_cudaData, int xRes, int yRes, Vec3d eyepoint, Vec3d up, Ve } } -// int index = 3*((sy-1) * xRes + sx); + int index = 3*((sy-1) * xRes + sx); d_cudaData[index + 0] = col[0]; d_cudaData[index + 1] = col[1]; d_cudaData[index + 2] = col[2]; + return; } From 85662c0905791429e65c2c79d0877a53649414e9 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 18:24:40 +0200 Subject: [PATCH 83/87] FIX true FPS and remove pseudo animation --- hscgv_uebung_5/ApplicationWindow.cpp | 31 +++------------------------- hscgv_uebung_5/ApplicationWindow.h | 6 ------ 2 files changed, 3 insertions(+), 34 deletions(-) diff --git a/hscgv_uebung_5/ApplicationWindow.cpp b/hscgv_uebung_5/ApplicationWindow.cpp index 74c7e72..b4d8e05 100644 --- a/hscgv_uebung_5/ApplicationWindow.cpp +++ b/hscgv_uebung_5/ApplicationWindow.cpp @@ -39,7 +39,6 @@ ApplicationWindow::ApplicationWindow() connect(ui.actionResetCamera, SIGNAL(triggered()), this, SIGNAL(resetCam())); connect(ui.actionResetLight, SIGNAL(triggered()), this, SIGNAL(resetLight())); - connect(ui.actionAnimate, SIGNAL(toggled(bool)), this, SLOT(animate(bool))); connect(ui.actionAntialiasing, SIGNAL(toggled(bool)), SLOT(antialiasing(bool))); connect(ui.actionGPU, SIGNAL(toggled(bool)), m_frame, SLOT(setRenderMode(bool))); @@ -65,10 +64,8 @@ ApplicationWindow::ApplicationWindow() // timer for continuous rendering m_trigger = new QTimer(this); - m_trigger->setInterval(1); // every ms connect(m_trigger, SIGNAL(timeout()), m_frame, SLOT(updateGL())); - - initState(); + m_trigger->start(); // load standard scene // if(qApp->argc() > 1) @@ -77,15 +74,6 @@ ApplicationWindow::ApplicationWindow() // loadFile("REF1.data"); } -void ApplicationWindow::initState() const -{ - // trigger toggles the action's state - - ui.actionAnimate->setChecked(true); - ui.actionAnimate->trigger(); -} - - // destroy the application window ApplicationWindow::~ApplicationWindow() { @@ -108,19 +96,14 @@ void ApplicationWindow::updateMessage(const QString& message) statusBar()->showMessage(message); } + void ApplicationWindow::updateFps() { static int oldFrames = 0; - m_secondsToDisplay--; int frames = m_frame->frameCounter() - oldFrames; oldFrames = m_frame->frameCounter(); - if(m_secondsToDisplay >= 0) - statusBar()->showMessage(m_message); - else if(frames > 0) - statusBar()->showMessage(QString::number(frames) + " f/s"); - else - statusBar()->showMessage("Use mouse buttons and wheel to move camera, hold shift to move light"); + statusBar()->showMessage(QString::number(frames) + " f/s"); } void ApplicationWindow::loadFile() @@ -144,14 +127,6 @@ void ApplicationWindow::loadFile(QString filename) emit openFile(filename); } -void ApplicationWindow::animate(bool on) -{ - if(on) - m_trigger->start(); - else - m_trigger->stop(); -} - void ApplicationWindow::antialiasing(bool on) { m_frame->m_antialiasing = on; diff --git a/hscgv_uebung_5/ApplicationWindow.h b/hscgv_uebung_5/ApplicationWindow.h index 8b3245d..302d7d4 100644 --- a/hscgv_uebung_5/ApplicationWindow.h +++ b/hscgv_uebung_5/ApplicationWindow.h @@ -63,9 +63,6 @@ class ApplicationWindow : public QMainWindow void renderMode(bool mode) const; protected: - //! initialize application state - void initState() const; - //! our main widget in the application is a OGLCanvas GLFrame *m_frame; @@ -93,9 +90,6 @@ class ApplicationWindow : public QMainWindow //! update FPS display void updateFps(); - //! start and stop render timer - void animate(bool on); - //! we want to have some nice output void antialiasing(bool on); From 38793bbd06147b7ce90f174df1c8977c91a24ea1 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Wed, 8 Jul 2015 18:25:06 +0200 Subject: [PATCH 84/87] IMPROVE handle m_data as a std::vector --- hscgv_uebung_5/GLFrame.cpp | 14 +++++++------- hscgv_uebung_5/GLFrame.h | 4 +++- 2 files changed, 10 insertions(+), 8 deletions(-) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index 8fa77b8..b84fade 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -58,7 +58,6 @@ GLFrame::GLFrame(ApplicationWindow *parent) GLFrame::~GLFrame() { glDeleteTextures(1, &m_texHandle); - free(m_data); cudaFree(m_cudaData); cudaDeviceReset(); } @@ -95,7 +94,7 @@ void GLFrame::loadTexture() glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER); - glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, m_width, m_height, 0, GL_RGB, GL_FLOAT, (GLvoid*)m_data); CHECKGL; + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB, m_width, m_height, 0, GL_RGB, GL_FLOAT, (GLvoid*)m_data.data()); CHECKGL; if (m_renderMode == GPU) { // register this texture with CUDA @@ -115,6 +114,8 @@ void GLFrame::resizeGL(int w, int h) m_width = w; m_height = h; + m_sizeTex = sizeof(float) * 3 * m_width * m_height; + m_data.resize(m_sizeTex); glViewport(0, 0, (GLint)w, (GLint)h); m_raytracingNeeded = true; updateGL(); @@ -562,17 +563,15 @@ void GLFrame::renderScene() if(!m_raytracingNeeded || !m_raytracer || !m_raytracer->m_isFileLoaded) return; // draw scene - int sizeTex = sizeof(float) * 3 * m_width * m_height; - m_data = (float*)malloc(sizeTex); switch(m_renderMode) { case CPU: - m_raytracer->render(m_data, m_width, m_height); + m_raytracer->render(m_data.data(), m_width, m_height); break; case GPU: - cudaMalloc(&m_cudaData,sizeTex); + cudaMalloc(&m_cudaData,m_sizeTex); m_raytracer->renderCuda(m_cudaData, m_width, m_height); - cudaMemcpy(m_data,m_cudaData, sizeTex, cudaMemcpyDeviceToHost); + cudaMemcpy(m_data.data(),m_cudaData, m_sizeTex, cudaMemcpyDeviceToHost); cudaFree(m_cudaData); break; } @@ -594,6 +593,7 @@ void GLFrame::loadScene(const QString &filename) updateGL(); } + int GLFrame::frameCounter() const { return m_frameCounter; diff --git a/hscgv_uebung_5/GLFrame.h b/hscgv_uebung_5/GLFrame.h index 8e603af..a4467dd 100644 --- a/hscgv_uebung_5/GLFrame.h +++ b/hscgv_uebung_5/GLFrame.h @@ -151,6 +151,8 @@ public slots: int m_width; //! viewport height int m_height; + //! size of data + int m_sizeTex; //! viewport aspect ratio (determined by window size) GLfloat m_aspect; //! field of view @@ -192,7 +194,7 @@ public slots: //! currently loaded scene Raytracer *m_raytracer; bool m_raytracingNeeded; - float* m_data; + std::vector m_data; float* m_cudaData; //! how our model is to be drawn From 24cb135e3df275ecf7b3dddddb8af9389984d9e0 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Thu, 9 Jul 2015 00:22:37 +0200 Subject: [PATCH 85/87] ADD some example scenes --- hscgv_uebung_5/data/a.data | 129 +++++++++++++++++++++++ hscgv_uebung_5/data/b.data | 127 +++++++++++++++++++++++ hscgv_uebung_5/data/c.data | 122 ++++++++++++++++++++++ hscgv_uebung_5/data/snooker3.data | 165 ++++++++++++++++++++++++++++++ 4 files changed, 543 insertions(+) create mode 100644 hscgv_uebung_5/data/a.data create mode 100644 hscgv_uebung_5/data/b.data create mode 100644 hscgv_uebung_5/data/c.data create mode 100644 hscgv_uebung_5/data/snooker3.data diff --git a/hscgv_uebung_5/data/a.data b/hscgv_uebung_5/data/a.data new file mode 100644 index 0000000..c7fe449 --- /dev/null +++ b/hscgv_uebung_5/data/a.data @@ -0,0 +1,129 @@ +#============================================================================== +# Programmierpraktikum Computergrafik +#------------------------------------------------------------------------------ +# Aufgabe 3: Lichtblick +#------------------------------------------------------------------------------ +# Created by Peter Kipfer +#============================================================================== + +#------------------------------------------------------------------------------ +# image parameters + +resolution 600 600 # image width and height in pixels +background 0.0 0.0 0.0 # background color black + +#------------------------------------------------------------------------------ +# viewing parameters + +eyepoint -100.0 150.0 150.0 # eye coordinates +lookat -15.0 25.0 -40.0 # point the camera is looking at +fovy 55 # angle of view +aspect 1.0 # width-height ratio +up 0.0 1.0 0.0 # this vector points upwards + +#------------------------------------------------------------------------------ +# global lighting + +ambience 0.0 0.0 0.0 # black + +#------------------------------------------------------------------------------ +# geometry + +# all coordinate values are defined within a right-handed system + +# quadrics are implicitly defined by the following equation +# Ax^2 + Bxy + Cxz + Dx + Ey^2 + Fyz + Gy + Hz^2 + Jz + K = 0 + +# a sphere is a special quadric: +# A = E = H = 1.0 ; B = C = F = 0.0 ; D = -2X , G = -2Y , J = -2Z +# K = X^2 + Y^2 + Z^2 - radius^2 +# with (X,Y,Z) being the midpoint of the sphere + +# the other way round, you get the sphere parameters from the implicit +# quadrics formula by +# X = -0.5 * D , Y = -0.5 * G , Z = -0.5 * J +# radius = SQR( (D^2 + G^2 + J^2) / 4 - K ) +# if and only if this surface really is a sphere + +# polyon vertices must be ordered in such a way, that they are ordered +# counterclockwise when looking at the polygon to give a 'outside' +# surface + +numsurfaces 7 + +# geometry description 1 : sphere at (0,0,-400) with radius 100 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 0.0 1.0 800.0 150000.0 + +# geometry description 2 : sphere at (200,50,-100) with radius 150 +object quadric 1.0 0.0 0.0 -400.0 1.0 0.0 -100.0 1.0 200.0 30000.0 + +# geometry description 3 : sphere at (0,-1200,-500) with radius 1000 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 2400.0 1.0 1000.0 690000.0 + +# geometry description 4 : sphere at (-100,25,-300) with radius 50 +object quadric 1.0 0.0 0.0 200.0 1.0 0.0 -50.0 1.0 600.0 98125.0 + +# geometry description 5 : sphere at (0,100,-250) with radius 25 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 -200.0 1.0 500.0 71875.0 + +# geometry description 6 : sphere at (0,100,-250) with radius 25 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 -250.0 1.0 100.0 -11875.0 + +# geometry description 7 : sphere at (0,100,-250) with radius 25 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 200.0 1.0 500.0 171875.0 + +#------------------------------------------------------------------------------ +# surface properties + +numproperties 3 + +# property 1 +ambient 0.2 0.0 0.0 # soft reddish filter +diffuse 0.7 0.7 0.7 # light silver-grey +specular 1.0 4 +mirror 0.8 + +# property 2 +ambient 0.0 0.2 0.0 # soft greenish filter +diffuse 0.5 0.5 0.5 # mid silver-grey +specular 1.0 4 +mirror 0.7 + +# property 3 +ambient 0.0 0.0 0.2 # soft blueish filter +diffuse 0.1 0.6 0.1 # green +specular 1.0 4 +mirror 0.2 + +#------------------------------------------------------------------------------ +# lighting + +numlights 3 + +# lightsource 1 +direction -0.1 1.0 0.1 +colour 1.0 0.5 0.1 + +# lightsource 2 +direction 0.1 0.2 1.0 +colour 0.2 0.2 0.2 + +# lightsource 3 +direction 1.0 2.0 3.0 +colour 0.0 0.5 1.0 + +#------------------------------------------------------------------------------ +# objects + +numobjects 7 + +# lookup table +object 1 3 +object 2 1 +object 3 2 +object 4 3 +object 5 1 +object 6 1 +object 7 1 + +#*** END OF FILE ************************************************************** diff --git a/hscgv_uebung_5/data/b.data b/hscgv_uebung_5/data/b.data new file mode 100644 index 0000000..01cc859 --- /dev/null +++ b/hscgv_uebung_5/data/b.data @@ -0,0 +1,127 @@ +#============================================================================== +# Programmierpraktikum Computergrafik +#------------------------------------------------------------------------------ +# Aufgabe 3: Lichtblick +#------------------------------------------------------------------------------ +# Created by Peter Kipfer +#============================================================================== + +#------------------------------------------------------------------------------ +# image parameters + +resolution 600 600 # image width and height in pixels +background 0.0 0.0 0.0 # background color black + +#------------------------------------------------------------------------------ +# viewing parameters + +eyepoint -100.0 150.0 150.0 # eye coordinates +lookat -15.0 25.0 -40.0 # point the camera is looking at +fovy 55 # angle of view +aspect 1.0 # width-height ratio +up 0.0 1.0 0.0 # this vector points upwards + +#------------------------------------------------------------------------------ +# global lighting + +ambience 0.0 0.0 0.0 # black + +#------------------------------------------------------------------------------ +# geometry + +# all coordinate values are defined within a right-handed system + +# quadrics are implicitly defined by the following equation +# Ax^2 + Bxy + Cxz + Dx + Ey^2 + Fyz + Gy + Hz^2 + Jz + K = 0 + +# a sphere is a special quadric: +# A = E = H = 1.0 ; B = C = F = 0.0 ; D = -2X , G = -2Y , J = -2Z +# K = X^2 + Y^2 + Z^2 - radius^2 +# with (X,Y,Z) being the midpoint of the sphere + +# the other way round, you get the sphere parameters from the implicit +# quadrics formula by +# X = -0.5 * D , Y = -0.5 * G , Z = -0.5 * J +# radius = SQR( (D^2 + G^2 + J^2) / 4 - K ) +# if and only if this surface really is a sphere + +# polyon vertices must be ordered in such a way, that they are ordered +# counterclockwise when looking at the polygon to give a 'outside' +# surface + +numsurfaces 6 + +# geometry description 1 : sphere at (0,0,-400) with radius 100 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 0.0 1.0 800.0 150000.0 + +# geometry description 2 : sphere at (200,50,-100) with radius 150 +object quadric 1.0 0.0 0.0 -400.0 1.0 0.0 -100.0 1.0 200.0 30000.0 + +# geometry description 3 : sphere at (0,-1200,-500) with radius 1000 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 2400.0 1.0 1000.0 690000.0 + +# geometry description 4 : sphere at (-100,25,-300) with radius 50 +object quadric 1.0 0.0 0.0 200.0 1.0 0.0 -50.0 1.0 600.0 98125.0 + +# geometry description 5 : sphere at (0,100,-250) with radius 25 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 -200.0 1.0 500.0 71875.0 + +# geometry description 6 : sphere at (0,100,-250) with radius 25 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 -250.0 1.0 100.0 -11875.0 + + + +#------------------------------------------------------------------------------ +# surface properties + +numproperties 3 + +# property 1 +ambient 0.2 0.0 0.0 # soft reddish filter +diffuse 0.7 0.7 0.7 # light silver-grey +specular 1.0 4 +mirror 0.8 + +# property 2 +ambient 0.0 0.2 0.0 # soft greenish filter +diffuse 0.5 0.5 0.5 # mid silver-grey +specular 1.0 4 +mirror 0.7 + +# property 3 +ambient 0.0 0.0 0.2 # soft blueish filter +diffuse 0.1 0.6 0.1 # green +specular 1.0 4 +mirror 0.2 + +#------------------------------------------------------------------------------ +# lighting + +numlights 3 + +# lightsource 1 +direction -0.1 1.0 0.1 +colour 1.0 0.5 0.1 + +# lightsource 2 +direction 0.1 0.2 1.0 +colour 0.2 0.2 0.2 + +# lightsource 3 +direction 1.0 2.0 3.0 +colour 0.0 0.5 1.0 + +#------------------------------------------------------------------------------ +# objects + +numobjects 6 + +# lookup table +object 1 3 +object 2 1 +object 3 2 +object 4 3 +object 5 1 +object 6 1 + +#*** END OF FILE ************************************************************** diff --git a/hscgv_uebung_5/data/c.data b/hscgv_uebung_5/data/c.data new file mode 100644 index 0000000..5f41127 --- /dev/null +++ b/hscgv_uebung_5/data/c.data @@ -0,0 +1,122 @@ +#============================================================================== +# Programmierpraktikum Computergrafik +#------------------------------------------------------------------------------ +# Aufgabe 3: Lichtblick +#------------------------------------------------------------------------------ +# Created by Peter Kipfer +#============================================================================== + +#------------------------------------------------------------------------------ +# image parameters + +resolution 600 600 # image width and height in pixels +background 0.0 0.0 0.0 # background color black + +#------------------------------------------------------------------------------ +# viewing parameters + +eyepoint -100.0 150.0 150.0 # eye coordinates +lookat -15.0 25.0 -40.0 # point the camera is looking at +fovy 55 # angle of view +aspect 1.0 # width-height ratio +up 0.0 1.0 0.0 # this vector points upwards + +#------------------------------------------------------------------------------ +# global lighting + +ambience 0.0 0.0 0.0 # black + +#------------------------------------------------------------------------------ +# geometry + +# all coordinate values are defined within a right-handed system + +# quadrics are implicitly defined by the following equation +# Ax^2 + Bxy + Cxz + Dx + Ey^2 + Fyz + Gy + Hz^2 + Jz + K = 0 + +# a sphere is a special quadric: +# A = E = H = 1.0 ; B = C = F = 0.0 ; D = -2X , G = -2Y , J = -2Z +# K = X^2 + Y^2 + Z^2 - radius^2 +# with (X,Y,Z) being the midpoint of the sphere + +# the other way round, you get the sphere parameters from the implicit +# quadrics formula by +# X = -0.5 * D , Y = -0.5 * G , Z = -0.5 * J +# radius = SQR( (D^2 + G^2 + J^2) / 4 - K ) +# if and only if this surface really is a sphere + +# polyon vertices must be ordered in such a way, that they are ordered +# counterclockwise when looking at the polygon to give a 'outside' +# surface + +numsurfaces 5 + +# geometry description 1 : sphere at (0,0,-400) with radius 100 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 0.0 1.0 800.0 150000.0 + +# geometry description 2 : sphere at (200,50,-100) with radius 150 +object quadric 1.0 0.0 0.0 -400.0 1.0 0.0 -100.0 1.0 200.0 30000.0 + +# geometry description 3 : sphere at (0,-1200,-500) with radius 1000 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 2400.0 1.0 1000.0 690000.0 + +# geometry description 4 : sphere at (-100,25,-300) with radius 50 +object quadric 1.0 0.0 0.0 200.0 1.0 0.0 -50.0 1.0 600.0 98125.0 + +# geometry description 5 : sphere at (0,100,-250) with radius 25 +object quadric 1.0 0.0 0.0 0.0 1.0 0.0 200.0 1.0 500.0 71875.0 + + +#------------------------------------------------------------------------------ +# surface properties + +numproperties 3 + +# property 1 +ambient 0.2 0.0 0.0 # soft reddish filter +diffuse 0.7 0.7 0.7 # light silver-grey +specular 1.0 4 +mirror 0.8 + +# property 2 +ambient 0.0 0.2 0.0 # soft greenish filter +diffuse 0.5 0.5 0.5 # mid silver-grey +specular 1.0 4 +mirror 0.7 + +# property 3 +ambient 0.0 0.0 0.2 # soft blueish filter +diffuse 0.1 0.6 0.1 # green +specular 1.0 4 +mirror 0.2 + +#------------------------------------------------------------------------------ +# lighting + +numlights 3 + +# lightsource 1 +direction -0.1 1.0 0.1 +colour 1.0 0.5 0.1 + +# lightsource 2 +direction 0.1 0.2 1.0 +colour 0.2 0.2 0.2 + +# lightsource 3 +direction 1.0 2.0 3.0 +colour 0.0 0.5 1.0 + +#------------------------------------------------------------------------------ +# objects + +numobjects 5 + +# lookup table +object 1 3 +object 2 1 +object 3 2 +object 4 3 +object 5 1 + +#*** END OF FILE ************************************************************** diff --git a/hscgv_uebung_5/data/snooker3.data b/hscgv_uebung_5/data/snooker3.data new file mode 100644 index 0000000..40d7a98 --- /dev/null +++ b/hscgv_uebung_5/data/snooker3.data @@ -0,0 +1,165 @@ +#============================================================================== +# Generated Ray Tracing File +#------------------------------------------------------------------------------ +# /home/stefan/Documents/3D/3D/res/snooker1.data +#------------------------------------------------------------------------------ +# Generated by Stefan Zellmann's Renderer +#============================================================================== + +#------------------------------------------------------------------------------ +# picture parameters + +resolution 1280 800 +background 0.12941177189350128173828125 0 0 + +#------------------------------------------------------------------------------ +# viewing parameters + +eyepoint 18.2749729156494140625 12.06964397430419921875 -37.9771728515625 +lookat -0.5 2 -5.5 +fovy 45 +up -0.1527545154094696044921875 0.965471804141998291015625 0.2110402286052703857421875 + +aspect 1.6 + +#------------------------------------------------------------------------------ +# global lighting + +ambience 0 0 0 + +#------------------------------------------------------------------------------ +# geometry + +numsurfaces 20 + +object quadric 1 0 0 16 1 0 -4 1 -13.8579998016357421875 112.0110396255340674542821943759918212890625 + +object quadric 1 0 0 8 1 0 -4 1 -13.8579998016357421875 64.0110396255340674542821943759918212890625 + +object quadric 1 0 0 0 1 0 -4 1 -13.8579998016357421875 48.0110396255340674542821943759918212890625 + +object quadric 1 0 0 -8 1 0 -4 1 -13.8579998016357421875 64.0110396255340674542821943759918212890625 + +object quadric 1 0 0 -16 1 0 -4 1 -13.8579998016357421875 112.0110396255340674542821943759918212890625 + +object quadric 1 0 0 12 1 0 -4 1 -6.927999973297119140625 47.99929590750122088138596154749393463134765625 + +object quadric 1 0 0 4 1 0 -4 1 -6.927999973297119140625 15.99929590750122088138596154749393463134765625 + +object quadric 1 0 0 -4 1 0 -4 1 -6.927999973297119140625 15.99929590750122088138596154749393463134765625 + +object quadric 1 0 0 -12 1 0 -4 1 -6.927999973297119140625 47.99929590750122088138596154749393463134765625 + +object quadric 1 0 0 8 1 0 -4 1 0 16 + +object quadric 1 0 0 0 1 0 -4 1 0 0 + +object quadric 1 0 0 -8 1 0 -4 1 0 16 + +object quadric 1 0 0 4 1 0 -4 1 6.927999973297119140625 15.99929590750122088138596154749393463134765625 + +object quadric 1 0 0 -4 1 0 -4 1 6.927999973297119140625 15.99929590750122088138596154749393463134765625 + +object quadric 1 0 0 0 1 0 -4 1 13.8579998016357421875 48.0110396255340674542821943759918212890625 + +object quadric 1 0 0 0 1 0 -4 1 -28 196 + +object quadric 1 0 0 0 1 0 -4 1 21.8579998016357421875 119.4430388320770362042821943759918212890625 + +object quadric 1 0 0 -22 1 0 -4 1 30 346 + +object quadric 1 0 0 15.3999996185302734375 1 0 -4 1 35 365.5399970626831418485380709171295166015625 + +object quadric 1 0 0 24 1 0 -4 1 45.59999847412109375 663.839965209961519576609134674072265625 + +object quadric 1 0 0 6 1 0 -4 1 50 634 + + +#------------------------------------------------------------------------------ +# surface properties + +numproperties 8 + +ambient 0.20000000298023223876953125 0.20000000298023223876953125 0.20000000298023223876953125 +diffuse 0.81176471710205078125 0.0470588244497776031494140625 0.0470588244497776031494140625 +specular 1 200 +mirror 0.100000001490116119384765625 + +ambient 0.20000000298023223876953125 0.20000000298023223876953125 0.20000000298023223876953125 +diffuse 0 0 0 +specular 1 200 +mirror 0.100000001490116119384765625 + +ambient 0.20000000298023223876953125 0.20000000298023223876953125 0.20000000298023223876953125 +diffuse 1 0.525490224361419677734375 0.525490224361419677734375 +specular 1 200 +mirror 0.100000001490116119384765625 + +ambient 0.20000000298023223876953125 0.20000000298023223876953125 0.20000000298023223876953125 +diffuse 0.01568627543747425079345703125 0 1 +specular 1 200 +mirror 0.100000001490116119384765625 + +ambient 0.20000000298023223876953125 0.20000000298023223876953125 0.20000000298023223876953125 +diffuse 0.20784313976764678955078125 0.113725490868091583251953125 0 +specular 1 200 +mirror 0.100000001490116119384765625 + +ambient 0.20000000298023223876953125 0.20000000298023223876953125 0.20000000298023223876953125 +diffuse 0.01960784383118152618408203125 0.3176470696926116943359375 0.007843137718737125396728515625 +specular 1 200 +mirror 0.100000001490116119384765625 + +ambient 0.20000000298023223876953125 0.20000000298023223876953125 0.20000000298023223876953125 +diffuse 0.93725490570068359375 0.93725490570068359375 0 +specular 1 200 +mirror 0.100000001490116119384765625 + +ambient 0.20000000298023223876953125 0.20000000298023223876953125 0.20000000298023223876953125 +diffuse 0.0 0.8 0.0 +specular 0.2 2 +mirror 0.75 + +#------------------------------------------------------------------------------ +# lighting + +numlights 2 + +direction 0 100 100 +colour 0.75686275959014892578125 0.75686275959014892578125 0.96078431606292724609375 + +direction 0 100 -10 +colour 0.941176474094390869140625 0.941176474094390869140625 0.745098054409027099609375 + +#------------------------------------------------------------------------------ +# objects + +numobjects 20 + +object 1 1 +object 2 1 +object 3 1 +object 4 1 +object 5 1 +object 6 1 +object 7 1 +object 8 1 +object 9 1 +object 10 1 +object 11 1 +object 12 1 +object 13 1 +object 14 1 +object 15 1 +object 16 2 +object 17 3 +object 18 4 +object 19 5 +object 20 6 + + +# Local Variables: +# mode: shell-script +# End: +#*** END OF FILE ************************************************************** + From fe5e5551eb5f363f93b69049c8aa26091c0a8650 Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Thu, 9 Jul 2015 00:22:51 +0200 Subject: [PATCH 86/87] FIX setup the aspect right on start --- hscgv_uebung_5/GLFrame.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/hscgv_uebung_5/GLFrame.cpp b/hscgv_uebung_5/GLFrame.cpp index b84fade..3b62bff 100644 --- a/hscgv_uebung_5/GLFrame.cpp +++ b/hscgv_uebung_5/GLFrame.cpp @@ -76,6 +76,7 @@ void GLFrame::initializeGL() glEnable(GL_DEPTH_TEST); // allocate a texture name glGenTextures( 1, &m_texHandle );CHECKGL; + g_scene.view.aspect = m_aspect; } void GLFrame::loadTexture() From 4ed7227dd1eb2a37c919a0d29b94a323cea8655c Mon Sep 17 00:00:00 2001 From: Jeronim Morina Date: Thu, 9 Jul 2015 00:23:37 +0200 Subject: [PATCH 87/87] release version --- hscgv_uebung_5/src.pro | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/hscgv_uebung_5/src.pro b/hscgv_uebung_5/src.pro index 3912e84..52cd536 100644 --- a/hscgv_uebung_5/src.pro +++ b/hscgv_uebung_5/src.pro @@ -22,8 +22,7 @@ QMAKE_YACCFLAGS = -y -d # comment these for a release version DEFINES *= TRACE VERBOSE -CONFIG += debug -CONFIG -= release +CONFIG += release FORMS += \