mirror/OpenSpace
1
0
Fork 0
mirror of https://github.com/OpenSpace/OpenSpace.git synced 2026-05-12 22:39:09 -05:00
Code Issues Packages Projects Releases Wiki Activity

A lot of intermediate changes

Browse Source 
- Compiles and runs on Linux and OS X
- Produces weird results on OS X
This commit is contained in:
Jonas Strandstedt
2014-03-27 13:38:58 -04:00
parent 0974fbaeda
commit 6d00b52240
21 changed files with 383 additions and 150 deletions
Download Patch File Download Diff File Expand all files Collapse all files
config/flareConfig.txt
+9 -9
View File
@@ -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
config/single.xml
+1 -1
View File
@@ -4,7 +4,7 @@
<Window fullScreen="false">
<!-- 16:9 aspect ratio -->
<Size x="640" y="360" />
<Pos x="1000" y="50.0" />
<Pos x="500" y="50.0" />
<Viewport>
<Pos x="0.0" y="0.0" />
<Size x="1.0" y="1.0" />
ext/ghoul
+1 -1
Submodule ext/ghoul updated: 58385daea3...13022c5f96
kernels/RaycasterTSP.cl
+78 -14
View File
@@ -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,
kernels/TSPTraversal.cl
+81 -20
View File
@@ -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);
src/flare/BrickManager.cpp
+8 -7
View File
@@ -2,7 +2,8 @@
* Author: Victor Sand (victor.sand@gmail.com)
*
*/
#include <GL/glew.h>
#include <ghoul/opengl/ghoul_gl.h>
#include <flare/BrickManager.h>
#include <flare/Texture3D.h>
#include <flare/Config.h>
@@ -108,9 +109,9 @@ bool BrickManager::ReadHeader() {
numValsTot_ = numBrickVals_*numBricksFrame_;
fseeko(file_, 0, SEEK_END);
off fileSize = ftello(file_);
off calcFileSize = static_cast<off>(numBricksTree_) *
static_cast<off>(brickSize_) + dataPos_;
off_t fileSize = ftello(file_);
off_t calcFileSize = static_cast<off_t>(numBricksTree_) *
static_cast<off_t>(brickSize_) + dataPos_;
if (fileSize != calcFileSize) {
ERROR("Sizes don't match");
@@ -366,9 +367,9 @@ bool BrickManager::DiskToPBO(BUFFER_INDEX _pboIndex) {
static_cast<std::ios::pos_type>(brickSize_);
*/
off offset = dataPos_ +
static_cast<off>(brickIndex) *
static_cast<off>(brickSize_);
off_t offset = dataPos_ +
static_cast<off_t>(brickIndex) *
static_cast<off_t>(brickSize_);
// Skip reading if all bricks in sequence is already in PBO
if (inPBO != sequence) {
src/flare/BrickManager.h
+1 -1
View File
@@ -112,7 +112,7 @@ private:
// C-style I/O
std::FILE *file_;
off dataPos_;
off_t dataPos_;
bool hasReadHeader_;
bool atlasInitialized_;
src/flare/CLManager.cpp
+18 -12
View File
@@ -2,14 +2,7 @@
* Author: Victor Sand (victor.sand@gmail.com)
*
*/
#include <GL/glew.h>
#ifndef _WIN32
#include <GL/glx.h>
#else
#include <Windows.h>
#include <WinUser.h>
#include <CL/cl_gl.h>
#endif
//#include <ghoul/opencl/ghoul_cl.hpp>
#include <flare/CLManager.h>
#include <flare/CLProgram.h>
#include <flare/TransferFunction.h>
@@ -18,6 +11,11 @@
#include <flare/Utils.h>
#include <sstream>
#include <ghoul/logging/logmanager.h>
namespace {
std::string _loggerCat = "CLManager";
}
using namespace osp;
CLManager::CLManager() {
@@ -34,7 +32,7 @@ CLManager::~CLManager() {
for (unsigned int i=0; i<NUM_QUEUE_INDICES; ++i) {
clReleaseCommandQueue(commandQueues_[i]);
}
clReleaseContext(context_);
//clReleaseContext(context_);
}
bool CLManager::InitPlatform() {
@@ -110,6 +108,7 @@ bool CLManager::InitDevices() {
}
bool CLManager::CreateContext() {
if (numPlatforms_ < 1) {
ERROR("Number of platforms < 1, can't create context");
return false;
@@ -119,7 +118,7 @@ bool CLManager::CreateContext() {
ERROR("Number of devices < 1, can't create context");
return false;
}
/*
// Create an OpenCL context with a reference to an OpenGL context
cl_context_properties contextProperties[] = {
#ifndef _WIN32
@@ -139,14 +138,21 @@ bool CLManager::CreateContext() {
// TODO Support more than one device?
context_ = clCreateContext(contextProperties, 1, &devices_[0], NULL,
NULL, &error_);
*/
bool success = _context.createContextFromGLContext();
if(!success)
LDEBUG("Could not create GL context");
return CheckSuccess(error_, "CreateContext()");
devices_[0] = _context.device();
return success;
//return CheckSuccess(error_, "CreateContext()");
}
bool CLManager::CreateCommandQueue() {
for (unsigned int i=0; i<NUM_QUEUE_INDICES; ++i) {
commandQueues_[i]=clCreateCommandQueue(context_, devices_[0], 0, &error_);
commandQueues_[i] = std::move(_context.createCommandQueue());
//commandQueues_[i]=clCreateCommandQueue(context_, devices_[0], 0, &error_);
if (!CheckSuccess(error_, "CreateCommandQueue()")) {
return false;
}
src/flare/CLManager.h
+11 -7
View File
@@ -6,11 +6,13 @@
#ifndef CL_MANAGER_H_
#define CL_MANAGER_H_
#ifndef _WIN32
#include <CL/cl.hpp>
#else
#include <CL/cl.h>
#endif
#include <ghoul/opencl/ghoul_cl.h>
#include <ghoul/opencl/clcontext.h>
#include <ghoul/opencl/clcommandqueue.h>
#include <ghoul/opencl/clprogram.h>
#include <ghoul/opencl/clkernel.h>
#include <ghoul/opencl/clutil.h>
#include <ghoul/opencl/clworksize.h>
#include <map>
#include <string>
#include <flare/KernelConstants.h>
@@ -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<std::string, CLProgram*> clPrograms_;
ghoul::opencl::CLContext _context;
};
src/flare/CLProgram.cpp
+83 -31
View File
@@ -2,16 +2,18 @@
* Author: Victor Sand (victor.sand@gmail.com)
*
*/
#include <GL/glew.h>
//#include <ghoul/opengl/ghoul_gl.h>
#include <flare/CLProgram.h>
#include <flare/CLManager.h>
#include <flare/TransferFunction.h>
#include <flare/Texture.h>
#include <flare/Utils.h>
#include <ghoul/logging/logmanager.h>
#include <ghoul/opencl/ghoul_cl.hpp>
#ifdef _WIN32
#include <CL/cl_gl.h>
#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;
}
src/flare/CLProgram.h
+1 -5
View File
@@ -6,11 +6,7 @@
#ifndef CL_PROGRAM_H_
#define CL_PROGRAM_H_
#ifndef _WIN32
#include <CL/cl.hpp>
#else
#include <CL/cl.h>
#endif
#include <ghoul/opencl/ghoul_cl.h>
#include <map>
#include <string>
#include <flare/KernelConstants.h>
src/flare/Raycaster.cpp
+22 -8
View File
@@ -3,10 +3,9 @@
*
*/
#include <ghoul/logging/logmanager.h>
#include <sgct.h>
#ifndef _WIN32
#include <GL/glx.h>
#endif
#include <ghoul/opengl/ghoul_gl.h>
#include <fstream>
#include <flare/Raycaster.h>
#include <flare/Texture2D.h>
@@ -25,6 +24,8 @@
#include <stdint.h>
#include <unistd.h> // 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<void*>(&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<int>(tsp_->NumValuesPerNode());
traversalConstants_.numOTNodes_ = static_cast<int>(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<void*>(&kernelConstants_),
src/flare/Raycaster.h
+5 -5
View File
@@ -13,11 +13,7 @@ TODO: Iteratively break away parts from it into other classes.
#include <string>
#include <vector>
#include <glm/glm.hpp>
#ifndef _WIN32
#include <CL/cl.hpp>
#else
#include <CL/cl.h>
#endif
#include <ghoul/opencl/ghoul_cl.h>
#include <flare/KernelConstants.h>
#include <boost/timer/timer.hpp>
#include <flare/TSP.h>
@@ -189,6 +185,10 @@ private:
// Timer and timer constants
boost::timer::cpu_timer timer_;
cl_mem cubeFrontCLmem;
cl_mem cubeBackCLmem;
};
}
src/flare/ShaderProgram.cpp
+1 -6
View File
@@ -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 <GL/glew.h>
#ifndef _WIN32
#include <GL/glfw3.h>
#else
#include <GL/glfw3.h>
#endif
#include <ghoul/opengl/ghoul_gl.h>
#include <flare/Utils.h>
#include <flare/ShaderProgram.h>
#include <flare/Texture2D.h>
src/flare/TSP.cpp
+5 -5
View File
@@ -165,7 +165,7 @@ bool TSP::CalculateSpatialError() {
for (unsigned int brick=0; brick<numTotalNodes_; ++brick) {
// Offset in file
off offset = dataPos_ + static_cast<off>(brick*numBrickVals*sizeof(float));
off_t offset = dataPos_ + static_cast<off_t>(brick*numBrickVals*sizeof(float));
fseeko(in, offset, SEEK_SET);
fread(reinterpret_cast<void*>(&buffer[0]),
@@ -212,7 +212,7 @@ bool TSP::CalculateSpatialError() {
lb!=coveredLeafBricks.end(); ++lb) {
// Read brick
off offset = dataPos_+static_cast<off>((*lb)*numBrickVals*sizeof(float));
off_t offset = dataPos_+static_cast<off_t>((*lb)*numBrickVals*sizeof(float));
fseeko(in, offset, SEEK_SET);
fread(reinterpret_cast<void*>(&buffer[0]),
static_cast<size_t>(numBrickVals)*sizeof(float), 1, in);
@@ -323,7 +323,7 @@ bool TSP::CalculateTemporalError() {
std::vector<float> voxelStdDevs(numBrickVals);
// Read the whole brick to fill the averages
off offset = dataPos_+static_cast<off>(brick*numBrickVals*sizeof(float));
off_t offset = dataPos_+static_cast<off_t>(brick*numBrickVals*sizeof(float));
fseeko(in, offset, SEEK_SET);
fread(reinterpret_cast<void*>(&voxelAverages[0]),
static_cast<size_t>(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<off>((*leaf*numBrickVals+voxel)*sizeof(float));
off_t sampleOffset = dataPos_ +
static_cast<off_t>((*leaf*numBrickVals+voxel)*sizeof(float));
fseeko(in, sampleOffset, SEEK_SET);
float sample;
fread(reinterpret_cast<void*>(&sample), sizeof(float), 1, in);
src/flare/TSP.h
+1 -1
View File
@@ -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
src/flare/Texture.cpp
+3 -1
View File
@@ -2,7 +2,9 @@
* Author: Victor Sand (victor.sand@gmail.com)
*
*/
#include <GL/glew.h>
#include <ghoul/opengl/ghoul_gl.h>
#include <flare/Texture.h>
#include <flare/Utils.h>
src/flare/Utils.cpp
+1 -1
View File
@@ -2,7 +2,7 @@
* Author: Victor Sand (victor.sand@gmail.com)
*
*/
#include <GL/glew.h>
#include <ghoul/opengl/ghoul_gl.h>
#include <flare/Utils.h>
unsigned int osp::CheckGLError(std::string _location) {
src/flare/flare.cpp
+49 -15
View File
@@ -17,6 +17,10 @@
#include <vector>
#include <iostream>
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) {
src/flare/flare.h
+1
View File
@@ -26,6 +26,7 @@
#define __FLARE_H__
#include <GL/glew.h>
#include <ghoul/logging/logmanager.h>
#include <sgct.h>
#include <flare/Animator.h>
#include <flare/Raycaster.h>
src/openspaceengine.cpp
+3
View File
@@ -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);