From 6d00b522404f568fc3e278f9a4231e2df34dca17 Mon Sep 17 00:00:00 2001 From: Jonas Strandstedt Date: Thu, 27 Mar 2014 13:38:58 -0400 Subject: [PATCH] A lot of intermediate changes - Compiles and runs on Linux and OS X - Produces weird results on OS X --- config/flareConfig.txt | 18 +++--- config/single.xml | 2 +- ext/ghoul | 2 +- kernels/RaycasterTSP.cl | 92 ++++++++++++++++++++++++----- kernels/TSPTraversal.cl | 101 +++++++++++++++++++++++++------- src/flare/BrickManager.cpp | 15 ++--- src/flare/BrickManager.h | 2 +- src/flare/CLManager.cpp | 30 ++++++---- src/flare/CLManager.h | 18 +++--- src/flare/CLProgram.cpp | 114 ++++++++++++++++++++++++++---------- src/flare/CLProgram.h | 6 +- src/flare/Raycaster.cpp | 30 +++++++--- src/flare/Raycaster.h | 10 ++-- src/flare/ShaderProgram.cpp | 7 +-- src/flare/TSP.cpp | 10 ++-- src/flare/TSP.h | 2 +- src/flare/Texture.cpp | 4 +- src/flare/Utils.cpp | 2 +- src/flare/flare.cpp | 64 +++++++++++++++----- src/flare/flare.h | 1 + src/openspaceengine.cpp | 3 + 21 files changed, 383 insertions(+), 150 deletions(-) diff --git a/config/flareConfig.txt b/config/flareConfig.txt index 2b85ad1057..8162252cb4 100644 --- a/config/flareConfig.txt +++ b/config/flareConfig.txt @@ -1,8 +1,8 @@ # Filenames # Don't change during runtime # (Transfer function values can be changed during runtime though) -tsp_filename /home/hhellteg/data/enlil_128_256_32.tsp -transferfunction_filename ../config/transferfunctions/fire.txt +tsp_filename /Users/jonasstrandstedt/Downloads/enlil_64_32_8.tsp +transferfunction_filename ../../../config/transferfunctions/fire.txt # Window dimensions # Don't change during runtime @@ -26,7 +26,7 @@ local_worksize_y 16 # Scaling division to make textures smaller # Saves OpenCL threads # (A factor of 2 results in half the number of threads per dimension etc) -texture_division_factor 1 +texture_division_factor 1 # Error tolerances # Use -1 for no tolerance @@ -48,12 +48,12 @@ raycaster_intensity 1.0 animator_refresh_interval 0.50 # Various paths -raycaster_kernel_filename ../kernels/RaycasterTSP.cl -tsp_traversal_kernel_filename ../kernels/TSPTraversal.cl -cube_shader_vert_filename ../shaders/cubeVert.glsl -cube_shader_frag_filename ../shaders/cubeFrag.glsl -quad_shader_vert_filename ../shaders/quadVert.glsl -quad_shader_frag_filename ../shaders/quadFrag.glsl +raycaster_kernel_filename ../../../kernels/RaycasterTSP.cl +tsp_traversal_kernel_filename ../../../kernels/TSPTraversal.cl +cube_shader_vert_filename ../../../shaders/cubeVert.glsl +cube_shader_frag_filename ../../../shaders/cubeFrag.glsl +quad_shader_vert_filename ../../../shaders/quadVert.glsl +quad_shader_frag_filename ../../../shaders/quadFrag.glsl # Model start_pitch -20.0 diff --git a/config/single.xml b/config/single.xml index 6b5fdeed67..8bd77227c2 100644 --- a/config/single.xml +++ b/config/single.xml @@ -4,7 +4,7 @@ - + diff --git a/ext/ghoul b/ext/ghoul index 58385daea3..13022c5f96 160000 --- a/ext/ghoul +++ b/ext/ghoul @@ -1 +1 @@ -Subproject commit 58385daea3eb6b55533c63f50b4d4d668fd3c143 +Subproject commit 13022c5f965401d8620a3caa679a61f32d9937fd diff --git a/kernels/RaycasterTSP.cl b/kernels/RaycasterTSP.cl index 1c76f579e6..de239b4c07 100644 --- a/kernels/RaycasterTSP.cl +++ b/kernels/RaycasterTSP.cl @@ -12,12 +12,76 @@ struct KernelConstants { int paddedBrickDim_; }; - +float3 CartesianToSpherical(float3 _cartesian); +float Lerp(float _v0, float _v1, float _d); +int LeftBST(int _bstNodeIndex, int _numValuesPerNode, int _numOTNodes, + bool _bstRoot, __global __read_only int *_tsp); +int RightBST(int _bstNodeIndex, int _numValuesPerNode, int _numOTNodes, + bool _bstRoot, __global __read_only int *_tsp); +int ChildNodeIndex(int _bstNodeIndex, + int *_timespanStart, + int *_timespanEnd, + int _timestep, + int _numValuesPerNode, + int _numOTNodes, + bool _bstRoot, + __global __read_only int *_tsp); +int BrickIndex(int _bstNodeIndex, int _numValuesPerNode, + __global __read_only int *_tsp); +bool IsBSTLeaf(int _bstNodeIndex, int _numValuesPerNode, + bool _bstRoot, __global __read_only int *_tsp); +bool IsOctreeLeaf(int _otNodeIndex, int _numValuesPerNode, + __global __read_only int *_tsp); +int OTChildIndex(int _otNodeIndex, int _numValuesPerNode, + int _child, + __global __read_only int *_tsp) ; +float TemporalError(int _bstNodeIndex, int _numValuesPerNode, + __global __read_only int *_tsp) ; +float SpatialError(int _bstNodeIndex, int _numValuesPerNode, + __global __read_only int *_tsp); +int3 BoxCoords(float3 _globalCoords, int _boxesPerAxis) ; +int3 AtlasBoxCoords(int _brickIndex, + __global __read_only int *_brickList); +float3 AtlasCoords(float3 _globalCoords, int _brickIndex, int _boxesPerAxis, + int _paddedBrickDim, int _level, + __global __read_only int *_brickList) ; +void SampleAtlas(float4 *_color, float3 _coords, int _brickIndex, + int _boxesPerAxis, int _paddedBrickDim, int _level, + const sampler_t _atlasSampler, + __read_only image3d_t _textureAtlas, + __read_only image2d_t _transferFunction, + const sampler_t _tfSampler, + __global __read_only int *_brickList); +bool TraverseBST(int _otNodeIndex, int *_brickIndex, + __constant __read_only struct KernelConstants *_constants, + __global __read_only int *_tsp, const int _timestep); +int EnclosingChild(float3 _P, float _boxMid, float3 _offset) ; +void UpdateOffset(float3 *_offset, float _boxDim, int _child) ; +float4 TraverseOctree(float3 _rayO, float3 _rayD, float _maxDist, + __read_only image3d_t _textureAtlas, + __constant struct KernelConstants *_constants, + __read_only image2d_t _transferFunction, + __global __read_only int *_tsp, + __global __read_only int *_brickList, + const int _timestep); +__kernel void RaycasterTSP(__read_only image2d_t _cubeFront, + __read_only image2d_t _cubeBack, + __write_only image2d_t _output, + __read_only image3d_t _textureAtlas, + __constant struct KernelConstants *_constants, + __read_only image2d_t _transferFunction, + //__global __read_only float *_transferFunction, + __global __read_only int *_tsp, + __global __read_only int *_brickList, + const int _timestep); + + + // Turn normalized [0..1] cartesian coordinates // to normalized spherical [0..1] coordinates float3 CartesianToSpherical(float3 _cartesian) { // Put cartesian in [-1..1] range first - _cartesian = (float3)(-1.0) + 2.0* _cartesian; + _cartesian = (float3)(-1.0) + 2.0f* _cartesian; float r = length(_cartesian); float theta, phi; if (r == 0.0) { @@ -26,7 +90,7 @@ float3 CartesianToSpherical(float3 _cartesian) { theta = acospi(_cartesian.z/r); phi = (M_PI + atan2(_cartesian.y, _cartesian.x)) / (2.0*M_PI); } - r = r / native_sqrt(3.0); + r = r / native_sqrt(3.0f); // Sampler ignores w component return (float3)(r, theta, phi); } @@ -218,8 +282,8 @@ float3 AtlasCoords(float3 _globalCoords, int _brickIndex, int _boxesPerAxis, void SampleAtlas(float4 *_color, float3 _coords, int _brickIndex, int _boxesPerAxis, int _paddedBrickDim, int _level, const sampler_t _atlasSampler, - __global __read_only image3d_t _textureAtlas, - __global __read_only image2d_t _transferFunction, + __read_only image3d_t _textureAtlas, + __read_only image2d_t _transferFunction, const sampler_t _tfSampler, __global __read_only int *_brickList) { @@ -235,7 +299,7 @@ void SampleAtlas(float4 *_color, float3 _coords, int _brickIndex, float sample = read_imagef(_textureAtlas, _atlasSampler, a4).x; // Composition float4 tf = read_imagef(_transferFunction, _tfSampler, (float2)(sample, 0.0)); - *_color += (1.0 - _color->w)*tf; + *_color += (1.0f - _color->w)*tf; } @@ -358,9 +422,9 @@ void UpdateOffset(float3 *_offset, float _boxDim, int _child) { } float4 TraverseOctree(float3 _rayO, float3 _rayD, float _maxDist, - __global __read_only image3d_t _textureAtlas, + __read_only image3d_t _textureAtlas, __constant struct KernelConstants *_constants, - __global __read_only image2d_t _transferFunction, + __read_only image2d_t _transferFunction, __global __read_only int *_tsp, __global __read_only int *_brickList, const int _timestep) { @@ -438,7 +502,7 @@ float4 TraverseOctree(float3 _rayO, float3 _rayD, float _maxDist, // Keep traversing the octree // Next box dimension - boxDim /= 2.0; + boxDim /= 2.0f; // Current mid point float boxMid = boxDim; @@ -469,12 +533,12 @@ float4 TraverseOctree(float3 _rayO, float3 _rayD, float _maxDist, } -__kernel void RaycasterTSP(__global __read_only image2d_t _cubeFront, - __global __read_only image2d_t _cubeBack, - __global __write_only image2d_t _output, - __global __read_only image3d_t _textureAtlas, +__kernel void RaycasterTSP(__read_only image2d_t _cubeFront, + __read_only image2d_t _cubeBack, + __write_only image2d_t _output, + __read_only image3d_t _textureAtlas, __constant struct KernelConstants *_constants, - __global __read_only image2d_t _transferFunction, + __read_only image2d_t _transferFunction, //__global __read_only float *_transferFunction, __global __read_only int *_tsp, __global __read_only int *_brickList, diff --git a/kernels/TSPTraversal.cl b/kernels/TSPTraversal.cl index 252d8ffef8..3152cc9300 100644 --- a/kernels/TSPTraversal.cl +++ b/kernels/TSPTraversal.cl @@ -9,11 +9,66 @@ struct TraversalConstants { float spatialTolerance_; }; +float3 CartesianToSpherical(float3 _cartesian); +int OctreeRootNodeIndex(); +int LeftBST(int _bstNodeIndex, int _numValuesPerNode, int _numOTNodes, + bool _bstRoot, __global __read_only int *_tsp); +int RightBST(int _bstNodeIndex, int _numValuesPerNode, int _numOTNodes, + bool _bstRoot, __global __read_only int *_tsp); +int ChildNodeIndex(int _bstNodeIndex, + int *_timespanStart, + int *_timespanEnd, + int _timestep, + int _numValuesPerNode, + int _numOTNodes, + bool _bstRoot, + __global __read_only int *_tsp); +int BrickIndex(int _bstNodeIndex, int _numValuesPerNode, + __global __read_only int *_tsp) ; +bool IsBSTLeaf(int _bstNodeIndex, int _numValuesPerNode, + bool _bstRoot, __global __read_only int *_tsp) ; +bool IsOctreeLeaf(int _otNodeIndex, int _numValuesPerNode, + __global __read_only int *_tsp) ; +int OTChildIndex(int _otNodeIndex, int _numValuesPerNode, + int _child, + __global __read_only int *_tsp); +void AddToList(int _brickIndex, + __global volatile int *_reqList); +float TemporalError(int _bstNodeIndex, int _numValuesPerNode, + __global __read_only int *_tsp) ; +float SpatialError(int _bstNodeIndex, int _numValuesPerNode, + __global __read_only int *_tsp); +bool TraverseBST(int _otNodeIndex, + int *_brickIndex, + int _timestep, + __constant struct TraversalConstants *_constants, + __global volatile int *_reqList, + __global __read_only int *_tsp); +int EnclosingChild(float3 _P, float _boxMid, float3 _offset); +void UpdateOffset(float3 *_offset, float _boxDim, int _child) ; +void TraverseOctree(float3 _rayO, + float3 _rayD, + float _maxDist, + __constant struct TraversalConstants *_constants, + __global volatile int *_reqList, + __global __read_only int *_tsp, + const int _timestep); +__kernel void TSPTraversal(__read_only image2d_t _cubeFront, + __read_only image2d_t _cubeBack, + __constant struct TraversalConstants *_constants, + __global __read_only int *_tsp, + __global int *_reqList, + const int _timestep) ; + + + + + // Turn normalized [0..1] cartesian coordinates // to normalized spherical [0..1] coordinates float3 CartesianToSpherical(float3 _cartesian) { // Put cartesian in [-1..1] range first - _cartesian = (float3)(-1.0) + 2.0* _cartesian; + _cartesian = (float3)(-1.0) + 2.0f* _cartesian; float r = length(_cartesian); float theta, phi; if (r == 0.0) { @@ -22,7 +77,7 @@ float3 CartesianToSpherical(float3 _cartesian) { theta = acospi(_cartesian.z/r); phi = (M_PI + atan2(_cartesian.y, _cartesian.x)) / (2.0*M_PI); } - r = r / native_sqrt(3.0); + r = r / native_sqrt(3.0f); // Sampler ignores w component return (float3)(r, theta, phi); } @@ -105,7 +160,7 @@ bool IsOctreeLeaf(int _otNodeIndex, int _numValuesPerNode, int OTChildIndex(int _otNodeIndex, int _numValuesPerNode, int _child, __global __read_only int *_tsp) { - int firstChild = _tsp[_otNodeIndex*_numValuesPerNode + 1]; + int firstChild = _tsp[_otNodeIndex*_numValuesPerNode+1]; return firstChild + _child; } @@ -275,21 +330,25 @@ void TraverseOctree(float3 _rayO, float stepsize = _constants->stepsize_; float3 P = _rayO; // Keep traversing until the sample point goes outside the unit cube - float traversed = 0.0; - while (traversed < _maxDist) { + float traversed = 0.0f; + + int max_iterations = 50; + int iterations = 0; + bool ok = stepsize > 0.0f && stepsize < fabs(_maxDist); + while (traversed < _maxDist && iterations < max_iterations) { // Reset traversal variables - float3 offset = (float3)(0.0); - float boxDim = 1.0; + float3 offset = (float3)(0.0f); + float boxDim = 1.0f; int child; // Init the octree node index to the root int otNodeIndex = OctreeRootNodeIndex(); - + // Start traversing octree // Rely on finding a leaf for loop termination + while (true) { - // See if the BST tree is good enough int brickIndex = 0; bool bstSuccess = TraverseBST(otNodeIndex, @@ -307,7 +366,7 @@ void TraverseOctree(float3 _rayO, AddToList(brickIndex, _reqList); // We are now done with this node, so go to next break; - + // If the BST lookup failed but the octree node is a leaf, // add the brick anyway (it is the BST leaf) } else if (IsOctreeLeaf(otNodeIndex, @@ -319,9 +378,9 @@ void TraverseOctree(float3 _rayO, // If the BST lookup failed and we can traverse the octree, // visit the child that encloses the point } else { - + // Next box dimension - boxDim = boxDim/2.0; + boxDim = boxDim/2.0f; // Current mid point float boxMid = boxDim; @@ -338,23 +397,24 @@ void TraverseOctree(float3 _rayO, UpdateOffset(&offset, boxDim, child); // Update node index to new node - int oldIndex = otNodeIndex; + //int oldIndex = otNodeIndex; otNodeIndex = OTChildIndex(otNodeIndex, _constants->numValuesPerNode_, child, _tsp); - - } - + } } // while traversing // Update - traversed += stepsize; - P += stepsize * _rayD; + iterations = iterations + 1; + traversed = traversed + stepsize; + P = P + stepsize * _rayD; } // while (traversed < maxDist) + + } -__kernel void TSPTraversal(__global __read_only image2d_t _cubeFront, - __global __read_only image2d_t _cubeBack, +__kernel void TSPTraversal(__read_only image2d_t _cubeFront, + __read_only image2d_t _cubeBack, __constant struct TraversalConstants *_constants, __global __read_only int *_tsp, __global int *_reqList, @@ -368,6 +428,7 @@ __kernel void TSPTraversal(__global __read_only image2d_t _cubeFront, // Read from color cube textures float4 cubeFrontColor = read_imagef(_cubeFront, sampler, intCoords); + if (length(cubeFrontColor.xyz) == 0.0) return; float4 cubeBackColor = read_imagef(_cubeBack, sampler, intCoords); diff --git a/src/flare/BrickManager.cpp b/src/flare/BrickManager.cpp index 623b417884..deea308ee6 100644 --- a/src/flare/BrickManager.cpp +++ b/src/flare/BrickManager.cpp @@ -2,7 +2,8 @@ * Author: Victor Sand (victor.sand@gmail.com) * */ -#include + +#include #include #include #include @@ -108,9 +109,9 @@ bool BrickManager::ReadHeader() { numValsTot_ = numBrickVals_*numBricksFrame_; fseeko(file_, 0, SEEK_END); - off fileSize = ftello(file_); - off calcFileSize = static_cast(numBricksTree_) * - static_cast(brickSize_) + dataPos_; + off_t fileSize = ftello(file_); + off_t calcFileSize = static_cast(numBricksTree_) * + static_cast(brickSize_) + dataPos_; if (fileSize != calcFileSize) { ERROR("Sizes don't match"); @@ -366,9 +367,9 @@ bool BrickManager::DiskToPBO(BUFFER_INDEX _pboIndex) { static_cast(brickSize_); */ - off offset = dataPos_ + - static_cast(brickIndex) * - static_cast(brickSize_); + off_t offset = dataPos_ + + static_cast(brickIndex) * + static_cast(brickSize_); // Skip reading if all bricks in sequence is already in PBO if (inPBO != sequence) { diff --git a/src/flare/BrickManager.h b/src/flare/BrickManager.h index 6f05f59c44..b8b6f0aa4d 100644 --- a/src/flare/BrickManager.h +++ b/src/flare/BrickManager.h @@ -112,7 +112,7 @@ private: // C-style I/O std::FILE *file_; - off dataPos_; + off_t dataPos_; bool hasReadHeader_; bool atlasInitialized_; diff --git a/src/flare/CLManager.cpp b/src/flare/CLManager.cpp index 06fe9431a7..709bb58bfe 100644 --- a/src/flare/CLManager.cpp +++ b/src/flare/CLManager.cpp @@ -2,14 +2,7 @@ * Author: Victor Sand (victor.sand@gmail.com) * */ -#include -#ifndef _WIN32 - #include -#else - #include - #include - #include -#endif +//#include #include #include #include @@ -18,6 +11,11 @@ #include #include +#include +namespace { + std::string _loggerCat = "CLManager"; +} + using namespace osp; CLManager::CLManager() { @@ -34,7 +32,7 @@ CLManager::~CLManager() { for (unsigned int i=0; i -#else -#include -#endif +#include +#include +#include +#include +#include +#include +#include #include #include #include @@ -120,11 +122,13 @@ private: char deviceName_[MAX_NAME_LENGTH]; char driverVersion_[MAX_NAME_LENGTH]; char platformVersion_[MAX_NAME_LENGTH]; - cl_context context_; - cl_command_queue commandQueues_[NUM_QUEUE_INDICES]; + //cl_context context_; + ghoul::opencl::CLCommandQueue commandQueues_[NUM_QUEUE_INDICES]; // Programs are mapped using strings std::map clPrograms_; + + ghoul::opencl::CLContext _context; }; diff --git a/src/flare/CLProgram.cpp b/src/flare/CLProgram.cpp index c1189f6d58..8fc6a4547a 100644 --- a/src/flare/CLProgram.cpp +++ b/src/flare/CLProgram.cpp @@ -2,16 +2,18 @@ * Author: Victor Sand (victor.sand@gmail.com) * */ -#include +//#include #include #include #include #include #include +#include +#include -#ifdef _WIN32 -#include -#endif +namespace { + std::string _loggerCat = "CLProgram"; +} using namespace osp; @@ -32,7 +34,7 @@ CLProgram::~CLProgram() { bool CLProgram::CreateProgram(std::string _fileName) { int numChars; char *source = ReadSource(_fileName, numChars); - program_ = clCreateProgramWithSource(clManager_->context_, 1, + program_ = clCreateProgramWithSource(clManager_->_context, 1, (const char**)&source, NULL, &error_); free(source); @@ -41,9 +43,11 @@ bool CLProgram::CreateProgram(std::string _fileName) { bool CLProgram::BuildProgram() { +std::string options = "-cl-opt-disable"; error_ = clBuildProgram(program_, (cl_uint)0, - NULL, NULL, NULL, NULL); + NULL, options.c_str(), NULL, NULL); if (error_ != CL_SUCCESS) { + LDEBUG("Could not build program " << getErrorString(error_)); // Print build log char * log; size_t logSize = 0; @@ -86,14 +90,26 @@ bool CLProgram::AddTexture(unsigned int _argNr, Texture *_texture, return false; break; case GL_TEXTURE_2D: - texture = clCreateFromGLTexture2D(clManager_->context_, _permissions, - GL_TEXTURE_2D, 0, - _texture->Handle(), &error_); - break; - case GL_TEXTURE_3D: - texture = clCreateFromGLTexture3D(clManager_->context_, _permissions, - GL_TEXTURE_3D, 0, - _texture->Handle(), &error_); +#ifdef CL_VERSION_1_2 + texture = clCreateFromGLTexture(clManager_->_context, _permissions, + GL_TEXTURE_2D, 0, + _texture->Handle(), &error_); +#else + texture = clCreateFromGLTexture2D(clManager_->_context, _permissions, + GL_TEXTURE_2D, 0, + _texture->Handle(), &error_); +#endif + break; + case GL_TEXTURE_3D: +#ifdef CL_VERSION_1_2 + texture = clCreateFromGLTexture(clManager_->_context, _permissions, + GL_TEXTURE_3D, 0, + _texture->Handle(), &error_); +#else + texture = clCreateFromGLTexture2D(clManager_->_context, _permissions, + GL_TEXTURE_3D, 0, + _texture->Handle(), &error_); +#endif break; default: ERROR("Unknown GL texture type"); @@ -125,14 +141,26 @@ bool CLProgram::AddTexture(unsigned int _argNr, Texture *_texture, return false; break; case GL_TEXTURE_2D: - _clTextureMem = clCreateFromGLTexture2D(clManager_->context_, _permissions, +#ifdef CL_VERSION_1_2 + _clTextureMem = clCreateFromGLTexture(clManager_->_context, _permissions, GL_TEXTURE_2D, 0, _texture->Handle(), &error_); +#else + _clTextureMem = clCreateFromGLTexture2D(clManager_->_context, _permissions, + GL_TEXTURE_2D, 0, + _texture->Handle(), &error_); +#endif break; case GL_TEXTURE_3D: - _clTextureMem = clCreateFromGLTexture3D(clManager_->context_, _permissions, - GL_TEXTURE_3D, 0, - _texture->Handle(), &error_); +#ifdef CL_VERSION_1_2 + _clTextureMem = clCreateFromGLTexture(clManager_->_context, _permissions, + GL_TEXTURE_3D, 0, + _texture->Handle(), &error_); +#else + _clTextureMem = clCreateFromGLTexture2D(clManager_->_context, _permissions, + GL_TEXTURE_3D, 0, + _texture->Handle(), &error_); +#endif break; default: ERROR("Unknown GL texture type"); @@ -176,7 +204,7 @@ bool CLProgram::AddBuffer(unsigned int _argNr, } MemArg ma; ma.size_ = sizeof(cl_mem); - ma.mem_ = clCreateBuffer(clManager_->context_, + ma.mem_ = clCreateBuffer(clManager_->_context, _allocMode | _permissions, (size_t)_sizeInBytes, _hostPtr, @@ -196,10 +224,15 @@ bool CLProgram::ReadBuffer(unsigned int _argNr, ERROR("ReadBuffer(): Could not find mem arg " << _argNr); return false; } - error_ = clEnqueueReadBuffer( - clManager_->commandQueues_[CLManager::EXECUTE], - memArgs_[(cl_uint)_argNr].mem_, _blocking, 0, _sizeInBytes, - _hostPtr, 0, NULL, NULL); + error_ = clEnqueueReadBuffer(clManager_->commandQueues_[CLManager::EXECUTE], + memArgs_[(cl_uint)_argNr].mem_, + _blocking, + 0, + _sizeInBytes, + _hostPtr, + 0, + NULL, + NULL); return clManager_->CheckSuccess(error_, "ReadBuffer"); } @@ -208,25 +241,36 @@ bool CLProgram::ReleaseBuffer(unsigned int _argNr) { ERROR("ReleaseBuffer(): Could not find mem arg " << _argNr); return false; } + //LDEBUG("Releasing memory"); error_ = clReleaseMemObject(memArgs_[(cl_uint)_argNr].mem_); return clManager_->CheckSuccess(error_, "ReleaseBuffer"); } bool CLProgram::PrepareProgram() { - +//ghoulFinishGL(); +/* +#ifdef __APPLE__ + + //glFlushRenderAPPLE(); + glFinish(); + //glFlush(); +#else +*/ // Let OpenCL take control of the shared GL textures for (auto it = OGLTextures_.begin(); it != OGLTextures_.end(); ++it) { - error_ = clEnqueueAcquireGLObjects( - clManager_->commandQueues_[CLManager::EXECUTE], 1, - &(it->second), 0, NULL, NULL); - + ghoul::opencl::CLCommandQueue q = clManager_->commandQueues_[CLManager::EXECUTE]; + cl_command_queue clq = q; + error_ = clEnqueueAcquireGLObjects(clq, 1, &(it->second), 0, NULL, NULL); + if (!clManager_->CheckSuccess(error_, "PrepareProgram")) { + LDEBUG("error: " << getErrorString(error_)); ERROR("Failed to enqueue GL object aqcuisition"); ERROR("Failing object: " << it->first); return false; } } +//#endif // Set up kernel arguments of non-shared items for (auto it=memArgs_.begin(); it!=memArgs_.end(); ++it) { @@ -276,14 +320,14 @@ bool CLProgram::LaunchProgram(unsigned int _gx, unsigned int _gy, } bool CLProgram::FinishProgram() { - + +//#ifndef __APPLE__ // Make sure kernel is done error_ = clFinish(clManager_->commandQueues_[CLManager::EXECUTE]); if (!clManager_->CheckSuccess(error_, "FinishProgram, clFinish")) { ERROR("Failed to finish program"); return false; } - // Release shared OGL objects for (auto it=OGLTextures_.begin(); it!=OGLTextures_.end(); ++it) { error_ = clEnqueueReleaseGLObjects( @@ -295,7 +339,15 @@ bool CLProgram::FinishProgram() { return false; } } - + /* +#else + error_ = clFinish(clManager_->commandQueues_[CLManager::EXECUTE]); + if (!clManager_->CheckSuccess(error_, "FinishProgram, clFinish")) { + ERROR("Failed to finish program"); + return false; + } +#endif +*/ return true; } diff --git a/src/flare/CLProgram.h b/src/flare/CLProgram.h index bc80eede1e..7297aa8e55 100644 --- a/src/flare/CLProgram.h +++ b/src/flare/CLProgram.h @@ -6,11 +6,7 @@ #ifndef CL_PROGRAM_H_ #define CL_PROGRAM_H_ -#ifndef _WIN32 -#include -#else -#include -#endif +#include #include #include #include diff --git a/src/flare/Raycaster.cpp b/src/flare/Raycaster.cpp index eb97bd4066..db94eabe01 100644 --- a/src/flare/Raycaster.cpp +++ b/src/flare/Raycaster.cpp @@ -3,10 +3,9 @@ * */ +#include #include -#ifndef _WIN32 - #include -#endif +#include #include #include #include @@ -25,6 +24,8 @@ #include #include // sync() + + using namespace osp; uint32_t ZOrder(uint16_t xPos, uint16_t yPos, uint16_t zPos) { @@ -348,7 +349,8 @@ bool Raycaster::InitPipeline() { // Allocate space for the brick request list // Use 0 as default value brickRequest_.resize(tsp_->NumTotalNodes(), 0); - + + glFinish(); // Run TSP traversal for timestep 0 if (!LaunchTSPTraversal(0)) { ERROR("InitPipeline() - failed to launch TSP traversal"); @@ -357,7 +359,7 @@ bool Raycaster::InitPipeline() { // Finish TSP traversal and read results into brick request if (!clManager_->FinishProgram("TSPTraversal")) return false; - + if (!clManager_->ReadBuffer("TSPTraversal", tspBrickListArg_, reinterpret_cast(&brickRequest_[0]), brickRequest_.size()*sizeof(int), @@ -670,14 +672,14 @@ bool Raycaster::InitCL() { } if (!clManager_->BuildProgram("TSPTraversal")) return false; if (!clManager_->CreateKernel("TSPTraversal")) return false; - cl_mem cubeFrontCLmem; + //cl_mem cubeFrontCLmem; if (!clManager_->AddTexture("TSPTraversal", tspCubeFrontArg_, cubeFrontTex_, CLManager::TEXTURE_2D, CLManager::READ_ONLY, cubeFrontCLmem)) { return false; } - cl_mem cubeBackCLmem; + //cl_mem cubeBackCLmem; if (!clManager_->AddTexture("TSPTraversal", tspCubeBackArg_, cubeBackTex_, CLManager::TEXTURE_2D, CLManager::READ_ONLY, cubeBackCLmem)) { @@ -690,6 +692,7 @@ bool Raycaster::InitCL() { CLManager::READ_ONLY)) return false; + LDEBUGC("RAYCASTER", config_->RaycasterKernelFilename()); // Raycaster part if (!clManager_->CreateProgram("RaycasterTSP", config_->RaycasterKernelFilename())) { @@ -769,7 +772,18 @@ bool Raycaster::UpdateKernelConstants() { static_cast(tsp_->NumValuesPerNode()); traversalConstants_.numOTNodes_ = static_cast(tsp_->NumOTNodes()); traversalConstants_.temporalTolerance_ = config_->TemporalErrorTolerance(); - traversalConstants_.spatialTolerance_ = config_->SpatialErrorTolerance(); + traversalConstants_.spatialTolerance_ = config_->SpatialErrorTolerance(); + + + std::string _loggerCat = "KOLLA KONSTANTER"; + LDEBUG("traversalConstants_.gridType_: " << traversalConstants_.gridType_); + LDEBUG("traversalConstants_.numOTNodes_: " << traversalConstants_.numOTNodes_); + LDEBUG("traversalConstants_.numTimesteps_: " << traversalConstants_.numTimesteps_); + LDEBUG("traversalConstants_.numValuesPerNode_: " << traversalConstants_.numValuesPerNode_); + LDEBUG("traversalConstants_.spatialTolerance_: " << traversalConstants_.spatialTolerance_); + LDEBUG("traversalConstants_.stepsize_: " << traversalConstants_.stepsize_); + LDEBUG("traversalConstants_.temporalTolerance_: " << traversalConstants_.temporalTolerance_); + if (!clManager_->AddBuffer("RaycasterTSP", constantsArg_, reinterpret_cast(&kernelConstants_), diff --git a/src/flare/Raycaster.h b/src/flare/Raycaster.h index ef86c18857..e739b1f2ab 100644 --- a/src/flare/Raycaster.h +++ b/src/flare/Raycaster.h @@ -13,11 +13,7 @@ TODO: Iteratively break away parts from it into other classes. #include #include #include -#ifndef _WIN32 -#include -#else -#include -#endif +#include #include #include #include @@ -189,6 +185,10 @@ private: // Timer and timer constants boost::timer::cpu_timer timer_; + + + cl_mem cubeFrontCLmem; + cl_mem cubeBackCLmem; }; } diff --git a/src/flare/ShaderProgram.cpp b/src/flare/ShaderProgram.cpp index 19b49d5b0b..bb7c7b2b0e 100644 --- a/src/flare/ShaderProgram.cpp +++ b/src/flare/ShaderProgram.cpp @@ -6,12 +6,7 @@ // TODO abstraction of shader binder, maybe a templated ShaderBinder class? // or string values service // possibly make a common UniformType class to handle matrices, ints, floats -#include -#ifndef _WIN32 - #include -#else - #include -#endif +#include #include #include #include diff --git a/src/flare/TSP.cpp b/src/flare/TSP.cpp index 954cee57e8..80587aabc3 100644 --- a/src/flare/TSP.cpp +++ b/src/flare/TSP.cpp @@ -165,7 +165,7 @@ bool TSP::CalculateSpatialError() { for (unsigned int brick=0; brick(brick*numBrickVals*sizeof(float)); + off_t offset = dataPos_ + static_cast(brick*numBrickVals*sizeof(float)); fseeko(in, offset, SEEK_SET); fread(reinterpret_cast(&buffer[0]), @@ -212,7 +212,7 @@ bool TSP::CalculateSpatialError() { lb!=coveredLeafBricks.end(); ++lb) { // Read brick - off offset = dataPos_+static_cast((*lb)*numBrickVals*sizeof(float)); + off_t offset = dataPos_+static_cast((*lb)*numBrickVals*sizeof(float)); fseeko(in, offset, SEEK_SET); fread(reinterpret_cast(&buffer[0]), static_cast(numBrickVals)*sizeof(float), 1, in); @@ -323,7 +323,7 @@ bool TSP::CalculateTemporalError() { std::vector voxelStdDevs(numBrickVals); // Read the whole brick to fill the averages - off offset = dataPos_+static_cast(brick*numBrickVals*sizeof(float)); + off_t offset = dataPos_+static_cast(brick*numBrickVals*sizeof(float)); fseeko(in, offset, SEEK_SET); fread(reinterpret_cast(&voxelAverages[0]), static_cast(numBrickVals)*sizeof(float), 1, in); @@ -349,8 +349,8 @@ bool TSP::CalculateTemporalError() { leaf != coveredBricks.end(); ++leaf) { // Sample the leaves at the corresponding voxel position - off sampleOffset = dataPos_ + - static_cast((*leaf*numBrickVals+voxel)*sizeof(float)); + off_t sampleOffset = dataPos_ + + static_cast((*leaf*numBrickVals+voxel)*sizeof(float)); fseeko(in, sampleOffset, SEEK_SET); float sample; fread(reinterpret_cast(&sample), sizeof(float), 1, in); diff --git a/src/flare/TSP.h b/src/flare/TSP.h index fbaced99fb..0e55448860 100644 --- a/src/flare/TSP.h +++ b/src/flare/TSP.h @@ -139,7 +139,7 @@ private: //std::ios::pos_type dataPos_; // Position of first data entry (after header) - off dataPos_; + off_t dataPos_; // Calculate weighted square distance between two RGBA colors // c2 should be an averaged or zero color diff --git a/src/flare/Texture.cpp b/src/flare/Texture.cpp index 9f849a3516..55154d980e 100644 --- a/src/flare/Texture.cpp +++ b/src/flare/Texture.cpp @@ -2,7 +2,9 @@ * Author: Victor Sand (victor.sand@gmail.com) * */ -#include + + +#include #include #include diff --git a/src/flare/Utils.cpp b/src/flare/Utils.cpp index fa7a71ffae..655aab852c 100644 --- a/src/flare/Utils.cpp +++ b/src/flare/Utils.cpp @@ -2,7 +2,7 @@ * Author: Victor Sand (victor.sand@gmail.com) * */ -#include +#include #include unsigned int osp::CheckGLError(std::string _location) { diff --git a/src/flare/flare.cpp b/src/flare/flare.cpp index ee52e1650e..948fcb3b1b 100644 --- a/src/flare/flare.cpp +++ b/src/flare/flare.cpp @@ -17,6 +17,10 @@ #include #include +namespace { + std::string _loggerCat = "Flare"; +} + namespace openspace { using namespace osp; @@ -39,6 +43,11 @@ Flare::~Flare() { delete _animator; } +void exit_msg(std::string m) { + LDEBUG(m); + exit(1); +} + void Flare::render() { // Go! // Reload config if flag is set @@ -55,6 +64,7 @@ void Flare::render() { _translateZ.getVal()); // Render if (!_raycaster->Render(_elapsedTime.getVal())) { + LDEBUG("!_raycaster->Render(_elapsedTime.getVal())"); exit(1); } @@ -90,7 +100,10 @@ void Flare::initNavigation() { void Flare::init() { // Start with reading a config file _config = Config::New(absPath("${BASE_PATH}/config/flareConfig.txt")); - if (!_config) exit(1); + if (!_config) { + LDEBUG("!_Config"); + exit(1); + } initNavigation(); _reloadFlag.setVal(false); @@ -115,26 +128,47 @@ void Flare::init() { // Create TSP structure from file TSP *tsp = TSP::New(_config); - if (!tsp->ReadHeader()) exit(1); + if (!tsp->ReadHeader()) { + LDEBUG("!tsp->ReadHeader()"); + exit(1); + } // Read cache if it exists, calculate otherwise if (tsp->ReadCache()) { INFO("\nUsing cached TSP file"); } else { INFO("\nNo cached TSP file found"); - if (!tsp->Construct()) exit(1); + if (!tsp->Construct()) { + LDEBUG("!tsp->Construct()"); + exit(1); + } if (_config->CalculateError() == 0) { INFO("Not calculating errors"); } else { - if (!tsp->CalculateSpatialError()) exit(1); - if (!tsp->CalculateTemporalError()) exit(1); - if (!tsp->WriteCache()) exit(1); + if (!tsp->CalculateSpatialError()) { + LDEBUG("!tsp->CalculateSpatialError()"); + exit(1); + } + if (!tsp->CalculateTemporalError()) { + LDEBUG("!tsp->CalculateTemporalError()"); + exit(1); + } + if (!tsp->WriteCache()) { + LDEBUG("!tsp->WriteCache()"); + exit(1); + } } } // Create brick manager and init (has to be done after init OpenGL!) BrickManager *brickManager= BrickManager::New(_config); - if (!brickManager->ReadHeader()) exit(1); - if (!brickManager->InitAtlas()) exit(1); + if (!brickManager->ReadHeader()) { + LDEBUG("!brickManager->ReadHeader()"); + exit(1); + } + if (!brickManager->InitAtlas()) { + LDEBUG("!brickManager->InitAtlas()"); + exit(1); + } // Create shaders for color cube and output textured quad ShaderProgram *cubeShaderProgram = ShaderProgram::New(); @@ -167,8 +201,8 @@ void Flare::init() { // Create transfer functions TransferFunction *transferFunction = TransferFunction::New(); transferFunction->SetInFilename(_config->TFFilename()); - if (!transferFunction->ReadFile()) exit(1); - if (!transferFunction->ConstructTexture()) exit(1); + if (!transferFunction->ReadFile()) exit_msg("!transferFunction->ReadFile()"); + if (!transferFunction->ConstructTexture()) exit_msg("!transferFunction->ConstructTexture()"); // Create animator @@ -183,15 +217,15 @@ void Flare::init() { _raycaster = Raycaster::New(_config); _raycaster->SetWinWidth(width); _raycaster->SetWinHeight(height); - if (!_raycaster->InitCube()) exit(1); - if (!_raycaster->InitQuad()) exit(1); + if (!_raycaster->InitCube()) exit_msg("!_raycaster->InitCube()"); + if (!_raycaster->InitQuad()) exit_msg("!_raycaster->InitQuad()"); _raycaster->SetBrickManager(brickManager); _raycaster->SetCubeFrontTexture(cubeFrontTex); _raycaster->SetCubeBackTexture(cubeBackTex); _raycaster->SetQuadTexture(quadTex); _raycaster->SetCubeShaderProgram(cubeShaderProgram); _raycaster->SetQuadShaderProgram(quadShaderProgram); - if (!_raycaster->InitFramebuffers()) exit(1); + if (!_raycaster->InitFramebuffers()) exit_msg("!_raycaster->InitFramebuffers()"); _raycaster->SetAnimator(_animator); _raycaster->AddTransferFunction(transferFunction); @@ -199,8 +233,8 @@ void Flare::init() { _raycaster->SetCLManager(clManager); _raycaster->SetTSP(tsp); - if (!_raycaster->InitCL()) exit(1); - if (!_raycaster->InitPipeline()) exit(1); + if (!_raycaster->InitCL()) exit_msg("!_raycaster->InitCL()"); + if (!_raycaster->InitPipeline()) exit_msg("!_raycaster->InitCL()"); } void Flare::keyboard(int key, int action) { diff --git a/src/flare/flare.h b/src/flare/flare.h index 1aec709396..af5f0ed300 100644 --- a/src/flare/flare.h +++ b/src/flare/flare.h @@ -26,6 +26,7 @@ #define __FLARE_H__ #include +#include #include #include #include diff --git a/src/openspaceengine.cpp b/src/openspaceengine.cpp index 381bc8e472..cd5d4f734b 100644 --- a/src/openspaceengine.cpp +++ b/src/openspaceengine.cpp @@ -171,6 +171,9 @@ bool OpenSpaceEngine::initialize() { ghoul::opencl::CLContext context; if(context.createContextFromGLContext()) { LDEBUG("Successfull CL/GL context creation"); + { + ghoul::opencl::CLContext tmp = context; + } ghoul::opencl::CLProgram prog = context.createProgram("${KERNELS}/test.cl"); prog.addDefinition("OFFSET", 3);