From a4a528375b270b6a5710b9cf66c07cc3c3f3f7d6 Mon Sep 17 00:00:00 2001 From: Jonas Strandstedt Date: Thu, 1 May 2014 14:17:48 -0400 Subject: [PATCH] Updated RenderableVolumeExpert to render properly - Added hack in VolumeRaycasteBox for SGCT Left/Right clearing - Fixed quad dimensions to depend on viewport size - Updated RenderableVolumeExpert to render correctly - Small changes to SceneGraph to make it easier to switch between volume rendering and Earth --- config/sgct/single.xml | 2 +- .../openspace/rendering/volumeraycasterbox.h | 15 ++- kernels/{ => helpers}/volume_helpers.cl | 6 +- kernels/helpers/volume_raycasting.cl | 104 ++++++++++++++++++ openspace-data | 2 +- openspace.cfg | 1 + src/engine/openspaceengine.cpp | 21 ++-- src/rendering/renderablevolumeexpert.cpp | 42 ++++--- src/rendering/volumeraycasterbox.cpp | 79 +++++++------ src/scenegraph/scenegraph.cpp | 2 +- 10 files changed, 197 insertions(+), 77 deletions(-) rename kernels/{ => helpers}/volume_helpers.cl (99%) create mode 100644 kernels/helpers/volume_raycasting.cl diff --git a/config/sgct/single.xml b/config/sgct/single.xml index 9f45468914..63abb4989c 100644 --- a/config/sgct/single.xml +++ b/config/sgct/single.xml @@ -1,7 +1,7 @@ - + diff --git a/include/openspace/rendering/volumeraycasterbox.h b/include/openspace/rendering/volumeraycasterbox.h index d32302f627..5792b0e0c3 100644 --- a/include/openspace/rendering/volumeraycasterbox.h +++ b/include/openspace/rendering/volumeraycasterbox.h @@ -25,13 +25,18 @@ #ifndef VOLUMERAYCASTERBOX_H_ #define VOLUMERAYCASTERBOX_H_ -#include -#include #include +// forward declare private objects namespace sgct_utils { class SGCTBox; } +namespace ghoul { + namespace opengl { + class FramebufferObject; + class ProgramObject; + } +} namespace openspace { @@ -40,7 +45,7 @@ public: VolumeRaycasterBox(); ~VolumeRaycasterBox(); bool initialize(); - void render(glm::mat4 MVP); + void render(const glm::mat4& MVP); ghoul::opengl::Texture* backFace(); ghoul::opengl::Texture* frontFace(); @@ -56,5 +61,5 @@ private: glm::size2_t _dimensions; }; -} /* namespace openspace */ -#endif /* VOLUMERAYCASTERBOX_H_ */ +} // namespace openspace +#endif // VOLUMERAYCASTERBOX_H_ diff --git a/kernels/volume_helpers.cl b/kernels/helpers/volume_helpers.cl similarity index 99% rename from kernels/volume_helpers.cl rename to kernels/helpers/volume_helpers.cl index 90ad1f4d9a..b3ba1d07ad 100644 --- a/kernels/volume_helpers.cl +++ b/kernels/helpers/volume_helpers.cl @@ -1,6 +1,5 @@ float3 CartesianToSpherical(float3 _cartesian); - float intensityNormalizer(float intensity, float iMin, float iMax); float3 CartesianToSpherical(float3 _cartesian) { @@ -23,4 +22,7 @@ float intensityNormalizer(float intensity, float iMin, float iMax) { float i = clamp(intensity, iMin, iMax); i = (i - iMin) / (iMax - iMin); return clamp(i, 0.0f, 1.0f); -} \ No newline at end of file +} + + + diff --git a/kernels/helpers/volume_raycasting.cl b/kernels/helpers/volume_raycasting.cl new file mode 100644 index 0000000000..e52c746b74 --- /dev/null +++ b/kernels/helpers/volume_raycasting.cl @@ -0,0 +1,104 @@ + +#define EARLY_RAY_TERMINATION_OPACITY 0.95 + +void raySetup(float3 first, float3 last, float3 dimension, float3* rayDirection, float* tIncr, float* tEnd); +bool earlyRayTermination(float4* color, float maxOpacity); + + +#define RC_DEFINE_TEXTUE_COORDINATES(coords) \ + int2 coords = (int2)(get_global_id(0), get_global_id(1)) + +#define RC_DEFINE_VOLUME3D_DIMENSIONS(dimension, volume) \ + float3 dimension; \ + { \ + int4 idim = get_image_dim(volume); \ + dimension = (float3)(idim.x,idim.y,idim.z); \ + } + + +//#define RC_DEFINE_VARIABLES(dimension, intCoords) + +/*** + * Calculates the direction of the ray and returns the number + * of steps and the direction. + ***/ +void raySetup(float3 first, float3 last, float3 dimension, float3* rayDirection, float* tIncr, float* tEnd) { + float samplingRate_ = 1.0f; + *rayDirection = last - first; + *tEnd = length(*rayDirection); + *rayDirection = normalize(*rayDirection); + *tIncr = 1.0/(samplingRate_*length((*rayDirection)*dimension)); + +} + +/*** + * Applies early ray termination. The current opacity is compared to + * the maximum opacity. In case it is greater, the opacity is set to + * 1.0 and true is returned, otherwise false is returned. + ***/ +bool earlyRayTermination(float4* color, float maxOpacity) { + if ((*color).a >= maxOpacity) { + (*color).a = 1.0f; + return true; + } else { + return false; + } +} + +#define RC_EARLY_RAY_TERMINATION(opacity, maxOpacity, finished) \ + finished = earlyRayTermination(&opacity, maxOpacity) + +#define RC_BEGIN_LOOP_FOR \ + for (int loop=0; !finished && loop tEnd); \ + RC_END_LOOP_BRACES +#else +#define RC_END_LOOP(result) \ + RC_EARLY_RAY_TERMINATION(result, EARLY_RAY_TERMINATION_OPACITY, finished); \ + t += tIncr; \ + finished = finished || (t > tEnd); \ + RC_END_LOOP_BRACES + +#endif + +/** +* In order to keep the shaders as free as possible from dealing +* with bricking and adaptive sampling, these defines can be placed +* before and after the compositing function calls to enable adaptive +* sampling when bricking is used. For normal datasets these defines +* will have no impact at all. +*/ +#ifdef ADAPTIVE_SAMPLING +#define RC_BEGIN_COMPOSITING \ + for (int i=0; i // sgct header has to be included before all others due to Windows header -#include "sgct.h" +#include #include #include @@ -34,8 +34,6 @@ #include #include - - #include #include #include @@ -45,7 +43,6 @@ #include #include #include - #include #include #include @@ -207,6 +204,12 @@ bool OpenSpaceEngine::isInitialized() { bool OpenSpaceEngine::initialize() { + // clear the screen so the user don't have to see old buffer contents from the graphics card + glClearColor(0,0,0,0); + glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); + GLFWwindow* win = sgct::Engine::instance()->getActiveWindowPtr()->getWindowHandle(); + glfwSwapBuffers(win); + // Register the filepaths from static function enables easy testing //registerFilePaths(); _context.createContextFromGLContext(); @@ -236,8 +239,6 @@ bool OpenSpaceEngine::initialize() { DeviceIdentifier::init(); DeviceIdentifier::ref().scanDevices(); _engine->_interactionHandler->connectDevices(); - - //_flare = new Flare(); return true; } @@ -274,8 +275,6 @@ void OpenSpaceEngine::preSynchronization() { _interactionHandler->update(dt); _interactionHandler->lockControls(); - - //_flare->preSync(); } } @@ -284,28 +283,24 @@ void OpenSpaceEngine::postSynchronizationPreDraw() { } void OpenSpaceEngine::render() { - //_flare->render(); _renderEngine->render(); } void OpenSpaceEngine::postDraw() { if (sgct::Engine::instance()->isMaster()) { _interactionHandler->unlockControls(); - //_flare->postDraw(); } } void OpenSpaceEngine::keyboardCallback(int key, int action) { if (sgct::Engine::instance()->isMaster()) { _interactionHandler->keyboardCallback(key, action); - //_flare->keyboard(key, action); } } void OpenSpaceEngine::mouseButtonCallback(int key, int action) { if (sgct::Engine::instance()->isMaster()) { _interactionHandler->mouseButtonCallback(key, action); - //_flare->mouse(key, action); } } @@ -318,11 +313,9 @@ void OpenSpaceEngine::mouseScrollWheelCallback(int pos) { } void OpenSpaceEngine::encode() { - //_flare->encode(); } void OpenSpaceEngine::decode() { - //_flare->decode(); } } // namespace openspace diff --git a/src/rendering/renderablevolumeexpert.cpp b/src/rendering/renderablevolumeexpert.cpp index 1b59e8938e..c9264789e5 100644 --- a/src/rendering/renderablevolumeexpert.cpp +++ b/src/rendering/renderablevolumeexpert.cpp @@ -298,6 +298,7 @@ bool RenderableVolumeExpert::initialize() { _colorBoxRenderer->initialize(); glm::size2_t dimensions = _colorBoxRenderer->dimensions(); + ghoul::opengl::Texture* backTexture = _colorBoxRenderer->backFace(); ghoul::opengl::Texture* frontTexture = _colorBoxRenderer->frontFace(); _output = new ghoul::opengl::Texture(glm::size3_t(dimensions[0],dimensions[1],1)); @@ -313,7 +314,7 @@ bool RenderableVolumeExpert::initialize() { // Compile kernels safeKernelCompilation(); - + // create work group size_t local_x = 32; size_t local_y = 32; while (local_x > 1) { @@ -341,32 +342,21 @@ void RenderableVolumeExpert::render(const Camera *camera, const psc &thisPositio if( ! _kernel.isValidKernel()) return; - glm::mat4 transform = camera->getViewProjectionMatrix(); - glm::mat4 camTransform = camera->getViewRotationMatrix(); psc camPos = camera->getPosition(); - -// psc relative = camPos - thisPosition; psc relative = thisPosition-camPos; double factor = pow(10.0,relative[3]); + glm::mat4 camTransform = camera->getViewRotationMatrix(); + glm::mat4 transform = camera->getViewProjectionMatrix(); transform = transform*camTransform; transform = glm::translate(transform, glm::vec3(relative[0]*factor, relative[1]*factor, relative[2]*factor)); transform = glm::scale(transform, _boxScaling); - /* -<<<<<<< HEAD - transform = glm::rotate(transform, time*speed, glm::vec3(0.0f, 1.0f, 0.0f)); - transform = glm::scale(transform, _boxScaling); -======= - transform = transform*camTransform; -// transform = glm::rotate(transform, speed*camera->getRotationAngle(), camera->getRotationAxis()); -// transform = glm::rotate(transform, time*speed, glm::vec3(0.0f, 1.0f, 0.0f)); - ->>>>>>> trackball -*/ _colorBoxRenderer->render(transform); _textureLock->lock(); _kernelLock->lock(); + + // tell opengl to finish everything before opencl takes ownerhip (uses) the textures glFinish(); // Aquire GL objects @@ -389,8 +379,12 @@ void RenderableVolumeExpert::render(const Camera *camera, const psc &thisPositio _quadProgram->activate(); glActiveTexture(GL_TEXTURE0); _output->bind(); - glClearColor(0.0f, 0.0f, 0.0f, 0.0f); - glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); + + // enable blending + glEnable(GL_BLEND); + glDisable(GL_DEPTH_TEST); + glBlendFunc(GL_ONE_MINUS_SRC_ALPHA, GL_SRC_ALPHA); + glBindVertexArray(_screenQuad); glDrawArrays(GL_TRIANGLES, 0, 6); glBindVertexArray(0); @@ -398,6 +392,11 @@ void RenderableVolumeExpert::render(const Camera *camera, const psc &thisPositio _kernelLock->unlock(); _textureLock->unlock(); + // disable blending + glDisable(GL_BLEND); + glEnable(GL_DEPTH_TEST); + glBlendFunc(GL_ONE, GL_ZERO); + _quadProgram->deactivate(); } @@ -451,7 +450,6 @@ void RenderableVolumeExpert::safeKernelCompilation() { if(argumentError) return; - tmpKernel.setArgument(0, &_clFrontTexture); tmpKernel.setArgument(1, &_clBackTexture); tmpKernel.setArgument(2, &_clOutput); @@ -496,24 +494,24 @@ void RenderableVolumeExpert::safeUpdateTexture(const ghoul::filesystem::File& fi // create the new texture ghoul::opengl::Texture* newTexture = loadTransferFunction(file.path()); - if(newTexture) { // upload the new texture and create a cl memory newTexture->uploadTexture(); cl_mem clNewTexture = _context.createTextureFromGLTexture(CL_MEM_READ_ONLY, *newTexture); + // check if opencl memory is unsuccessfull if(clNewTexture == 0) { delete newTexture; return; } - // everything is ok, critical point to replace current texture pointers + // everything seems ok, critical point to replace current texture pointers _textureLock->lock(); // deallocate current texture - delete _transferFunctions.at(fileID); clReleaseMemObject(_clTransferFunctions.at(fileID)); + delete _transferFunctions.at(fileID); // set the new texture _transferFunctions.at(fileID) = newTexture; diff --git a/src/rendering/volumeraycasterbox.cpp b/src/rendering/volumeraycasterbox.cpp index 59eac9b173..3e27e0d95d 100644 --- a/src/rendering/volumeraycasterbox.cpp +++ b/src/rendering/volumeraycasterbox.cpp @@ -27,6 +27,8 @@ #include #include +#include +#include #include #include @@ -39,29 +41,25 @@ using namespace ghoul::opengl; namespace openspace { -VolumeRaycasterBox::VolumeRaycasterBox() { +VolumeRaycasterBox::VolumeRaycasterBox(): _fbo(nullptr), _backTexture(nullptr), + _frontTexture(nullptr), _boxProgram(nullptr), _boundingBox(nullptr) { } -VolumeRaycasterBox::~VolumeRaycasterBox() {} +VolumeRaycasterBox::~VolumeRaycasterBox() { + if(_boundingBox) + delete _boundingBox; + + if(_fbo) { + _fbo->detachAll(); // maybe not needed + delete _fbo; + } +} bool VolumeRaycasterBox::initialize() { _boundingBox = new sgct_utils::SGCTBox(1.0f, sgct_utils::SGCTBox::Regular); // ------ SETUP SHADER ----------------- - _boxProgram = new ProgramObject("RaycastBoxProgram"); - /* - ShaderObject* vertexShader = new ShaderObject(ShaderObject::ShaderTypeVertex, - absPath("${BASE_PATH}/shaders/exitpoints.vert")); - ShaderObject* fragmentShader = new ShaderObject(ShaderObject::ShaderTypeFragment, - absPath("${BASE_PATH}/shaders/exitpoints.frag")); - - - _boxProgram->attachObject(vertexShader); - _boxProgram->attachObject(fragmentShader); - _boxProgram->compileShaderObjects(); - _boxProgram->linkProgramObject(); - */ OsEng.configurationManager().getValue("RaycastProgram", _boxProgram); _MVPLocation = _boxProgram->uniformLocation("modelViewProjection"); @@ -69,49 +67,68 @@ bool VolumeRaycasterBox::initialize() { _fbo = new FramebufferObject(); _fbo->activate(); + // changed from getActiveXResolution to getCurrentViewportPixelCoords because + // if there are more viewports in the same screen. + //size_t x = sgct::Engine::instance()->getActiveXResolution(); + //size_t y = sgct::Engine::instance()->getActiveYResolution(); + int x1, xSize, y1, ySize; + sgct::Engine::instance()->getActiveWindowPtr()->getCurrentViewportPixelCoords(x1, y1, xSize, ySize); + size_t x = xSize; + size_t y = ySize; - size_t x = sgct::Engine::instance()->getActiveXResolution(); - size_t y = sgct::Engine::instance()->getActiveYResolution(); _dimensions = glm::size2_t(x, y); - _backTexture = new ghoul::opengl::Texture(glm::size3_t(x,y,1)); - _frontTexture = new ghoul::opengl::Texture(glm::size3_t(x,y,1)); + _backTexture = new Texture(glm::size3_t(x,y,1)); + _frontTexture = new Texture(glm::size3_t(x,y,1)); _backTexture->uploadTexture(); _frontTexture->uploadTexture(); _fbo->attachTexture(_backTexture, GL_COLOR_ATTACHMENT0); _fbo->attachTexture(_frontTexture, GL_COLOR_ATTACHMENT1); - _fbo->deactivate(); return true; } -void VolumeRaycasterBox::render(glm::mat4 MVP) { - GLuint activeFBO = ghoul::opengl::FramebufferObject::getActiveObject(); // Save SGCTs main FBO +void VolumeRaycasterBox::render(const glm::mat4& MVP) { + GLuint activeFBO = FramebufferObject::getActiveObject(); // Save SGCTs main FBO _fbo->activate(); _boxProgram->activate(); _boxProgram->setUniform(_MVPLocation, MVP); + + sgct_core::Frustum::FrustumMode mode = sgct::Engine::instance()-> + getActiveWindowPtr()-> + getCurrentViewport()-> + getEye(); + + // oh god why..? + if(mode == sgct_core::Frustum::FrustumMode::Mono || + mode == sgct_core::Frustum::FrustumMode::StereoLeftEye) { + glDrawBuffer(GL_COLOR_ATTACHMENT0); + glClearColor(0.0, 0.0, 0.0, 0.0); + glClear(GL_COLOR_BUFFER_BIT); + glDrawBuffer(GL_COLOR_ATTACHMENT1); + glClearColor(0.0, 0.0, 0.0, 0.0); + glClear(GL_COLOR_BUFFER_BIT); + + } + // make sure GL_CULL_FACE is enabled (it should be) + glEnable(GL_CULL_FACE); + // Draw backface glDrawBuffer(GL_COLOR_ATTACHMENT0); - glClearColor(0.2f, 0.2f, 0.2f, 0); - glClear(GL_COLOR_BUFFER_BIT); - glEnable(GL_CULL_FACE); glCullFace(GL_FRONT); _boundingBox->draw(); - glDisable(GL_CULL_FACE); - // Draw frontface + // Draw frontface (now the normal cull face is is set) glDrawBuffer(GL_COLOR_ATTACHMENT1); - glClear(GL_COLOR_BUFFER_BIT); - glClearColor(0.2f, 0.2f, 0.2f, 0); - glEnable(GL_CULL_FACE); glCullFace(GL_BACK); _boundingBox->draw(); - glDisable(GL_CULL_FACE); _boxProgram->deactivate(); _fbo->deactivate(); + + // rebind the previous FBO glBindFramebuffer(GL_FRAMEBUFFER, activeFBO); } diff --git a/src/scenegraph/scenegraph.cpp b/src/scenegraph/scenegraph.cpp index d9b7ee98b4..fa097df24d 100644 --- a/src/scenegraph/scenegraph.cpp +++ b/src/scenegraph/scenegraph.cpp @@ -171,7 +171,7 @@ bool SceneGraph::initialize() { // TODO: Set scaling dependent on the position and distance // set position for camera psc cameraPosition = positionNode->getPosition(); - cameraPosition += psc(0.0,0.0,1.0,0.0); + cameraPosition += psc(0.0,0.0,2.0,0.0); c->setPosition(cameraPosition); c->setCameraDirection(glm::vec3(0,0,-1)); c->setScaling(glm::vec2(1.0,0.0));