Starting from 2021-07-01, all LRZ GitLab users will be required to explicitly accept the GitLab Terms of Service. Please see the detailed information at https://doku.lrz.de/display/PUBLIC/GitLab and make sure that your projects conform to the requirements.

Commit 576bad56 authored by schultezub's avatar schultezub
Browse files

* fixed KissCL deinitialization

 * fixed Program::getBuildLog()
 * CLRaycaster renders images now

git-svn-id: https://camplinux.in.tum.de/svn/campvis/trunk@246 bb408c1c-ae56-11e1-83d9-df6b3e0c105e
parent ab95ee8a
......@@ -144,10 +144,6 @@ namespace TUMVis {
// Deinit everything OpenGL related using the local context.
tgt::GLContextScopedLock lock(_localContext->getContext());
if (_useOpenCL) {
kisscl::CLRuntime::deinit();
}
// Deinit pipeline first
for (std::vector<AbstractPipeline*>::iterator it = _pipelines.begin(); it != _pipelines.end(); ++it) {
(*it)->deinit();
......@@ -160,6 +156,10 @@ namespace TUMVis {
// deinit OpenGL and tgt
tgt::deinitGL();
if (_useOpenCL) {
kisscl::CLRuntime::deinit();
}
}
tgt::QtContextManager::deinit();
......
......@@ -230,13 +230,13 @@ namespace TUMVis {
cl_channel_type WeaklyTypedPointer::getClChannelType() const {
switch (_baseType) {
case WeaklyTypedPointer::UINT8:
return CL_UNSIGNED_INT8;
return CL_UNORM_INT8;
case WeaklyTypedPointer::INT8:
return CL_SIGNED_INT8;
return CL_SNORM_INT8;
case WeaklyTypedPointer::UINT16:
return CL_UNSIGNED_INT16;
return CL_UNORM_INT16;
case WeaklyTypedPointer::INT16:
return CL_SIGNED_INT16;
return CL_SNORM_INT16;
case WeaklyTypedPointer::UINT32:
return CL_UNSIGNED_INT32;
case WeaklyTypedPointer::INT32:
......
......@@ -355,7 +355,7 @@ namespace kisscl {
* Gets the handle to internal OpenCl object.
* \return _id
*/
cl_type getId() const {
const cl_type& getId() const {
return _id;
}
......
......@@ -106,7 +106,7 @@ namespace kisscl {
switch(texFormat) {
case GL_ALPHA:
case GL_LUMINANCE:
imageFormat.image_channel_order = CL_A;
imageFormat.image_channel_order = CL_INTENSITY;
break;
case GL_LUMINANCE_ALPHA:
imageFormat.image_channel_order = CL_RA;
......
......@@ -95,7 +95,7 @@ namespace kisscl {
cl_int err = LCL_ERROR(clBuildProgram(_id, 0, 0, _buildOptions.c_str(), 0, 0));
if (err != CL_SUCCESS) {
for(size_t i = 0; i < _context->getDevices().size(); ++i)
LERROR(getBuildLog(_context->getDevices()[i]));
LERROR("Build log: " << getBuildLog(_context->getDevices()[i]));
}
}
else {
......@@ -141,7 +141,7 @@ namespace kisscl {
}
std::string Program::getBuildLog(const Device* device) const {
return getBuildInfo<std::string>(device, CL_PROGRAM_BUILD_OPTIONS);
return getBuildInfo<std::string>(device, CL_PROGRAM_BUILD_LOG);
}
//template specialization for strings:
......
......@@ -71,6 +71,7 @@ namespace TUMVis {
_camera.addSharedProperty(&_eepGenerator._camera);
_camera.addSharedProperty(&_drrraycater._camera);
_camera.addSharedProperty(&_simpleRaycaster._camera);
_camera.addSharedProperty(&_clRaycaster._camera);
_imageReader._url.setValue("D:\\Medical Data\\smallHeart.mhd");
_imageReader._targetImageID.setValue("reader.output");
......
......@@ -45,7 +45,7 @@ __constant float SAMPLING_BASE_INTERVAL_RCP = 200.0;
/**
* Makes a simple raycast through the volume for entry to exit point with minimal diffuse shading.
*/
float4 simpleRaycast(__global read_only image3d_t volumeTex, __global read_only image2d_t tfData, const float4 entryPoint, const float4 exitPoint, float* depth, const float stepSize) {
float4 simpleRaycast(__global read_only image3d_t volumeTex, __global read_only image2d_t tfData, const float4 entryPoint, const float4 exitPoint, float* depth, const float stepSize, float tfLowerBound, float tfUpperBound) {
// result color
float4 result = (float4)(0.0, 0.0, 0.0, 0.0);
......@@ -61,14 +61,19 @@ float4 simpleRaycast(__global read_only image3d_t volumeTex, __global read_only
//calculate sample position and get corresponding voxel
float4 sample = entryPoint + t * direction;
float4 color = read_imagef(tfData, smpNorm, (float2)(read_imagef(volumeTex, smpNorm, sample).x, 0.0));
float intensity = read_imagef(volumeTex, smpNorm, sample).w;
// apply tf intensity domain mapping:
intensity = (intensity - tfLowerBound) / (tfUpperBound - tfLowerBound);
float4 color = read_imagef(tfData, smpNorm, (float2)(intensity, 0.0));
// apply opacity correction to accomodate for variable sampling intervals
color.w = 1.0 - pow(1.0 - color.w, stepSize * SAMPLING_BASE_INTERVAL_RCP);
// Add a little shading. calcGradient is declared in mod_gradients.cl
float4 norm = normalize(calcGradient(sample, volumeTex));
color *= fabs(dot(norm, direction));
//float4 norm = normalize(calcGradient(sample, volumeTex));
//color *= fabs(dot(norm, direction));
//calculate ray integral
result.xyz = result.xyz + (1.0 - result.w) * color.w * color.xyz;
......@@ -98,7 +103,9 @@ __kernel void clraycaster(__global read_only image3d_t volumeTex,
__global read_only image2d_t entryTexCol,
__global read_only image2d_t exitTexCol,
__global write_only image2d_t outCol,
float stepSize
float stepSize,
float tfLowerBound,
float tfUpperBound
)
{
//output image pixel coordinates
......@@ -113,7 +120,7 @@ __kernel void clraycaster(__global read_only image3d_t volumeTex,
float4 exit = read_imagef(exitTexCol, smpNorm, targetNorm);
if( entry.x != exit.x || entry.y != exit.y || entry.z != exit.z )
color = simpleRaycast(volumeTex, tfData, entry, exit, &depth, stepSize);
color = simpleRaycast(volumeTex, tfData, entry, exit, &depth, stepSize, tfLowerBound, tfUpperBound);
else
color = (float4)(0.0);
......
......@@ -93,9 +93,8 @@ namespace TUMVis {
_clContext = CLRtm.createGlSharingContext();
_clProgram = CLRtm.loadProgram(_clContext, "modules/vis/clraycaster.cl");
_clProgram->setBuildOptions(" -cl-fast-relaxed-math -cl-mad-enable");
_clProgram->build();
std::string log = _clProgram->getBuildLog(_clContext->getDevices().front());
LDEBUG(log);
}
void CLRaycaster::deinit() {
......@@ -132,20 +131,28 @@ namespace TUMVis {
_texExitPointsColor = new kisscl::SharedTexture(_clContext, CL_MEM_READ_ONLY, exitPoints->getColorTexture());
delete _texOutColor;
ImageDataRenderTarget* rt = new ImageDataRenderTarget(tgt::svec3(_renderTargetSize.getValue(), 1));
_texOutColor = new kisscl::SharedTexture(_clContext, CL_MEM_READ_ONLY, rt->getColorTexture());
_texOutColor = new kisscl::SharedTexture(_clContext, CL_MEM_WRITE_ONLY, rt->getColorTexture());
// prepare kernel and stuff command queue
kisscl::CommandQueue* cq = CLRtm.getCommandQueue(_clContext);
kisscl::Kernel* kernel = _clProgram->getKernel("foobar");
kisscl::Kernel* kernel = _clProgram->getKernel("clraycaster");
if (kernel != 0) {
//kernel->setMemoryArgument(0, _imgVolume);
//kernel->setMemoryArgument(1, _imgTf);
kernel->setMemoryArgument(0, _texEntryPointsColor);
kernel->setMemoryArgument(1, _texExitPointsColor);
kernel->setMemoryArgument(2, _texOutColor);
//kernel->setArgument(5, _samplingStepSize.getValue());
rt->activate();
rt->deactivate();
LGL_ERROR;
glFinish();
kernel->setMemoryArgument(0, _imgVolume);
kernel->setMemoryArgument(1, _imgTf);
kernel->setMemoryArgument(2, _texEntryPointsColor);
kernel->setMemoryArgument(3, _texExitPointsColor);
kernel->setMemoryArgument(4, _texOutColor);
kernel->setArgument(5, _samplingStepSize.getValue());
kernel->setArgument(6, _transferFunction.getTF()->getIntensityDomain().x);
kernel->setArgument(7, _transferFunction.getTF()->getIntensityDomain().y);
cq->enqueueAcquireGLObject(_texEntryPointsColor);
cq->enqueueAcquireGLObject(_texExitPointsColor);
......@@ -191,6 +198,7 @@ namespace TUMVis {
}*/
LGL_ERROR;
data.addData(_targetImageID.getValue(), rt);
}
else {
LERROR("Input image must have dimensionality of 3.");
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment