View previous topic :: View next topic |
|
Author |
Message |
serkan.ergun Newbie
Joined: 01 Feb 2011 Posts: 3
|
Posted: Tue Feb 01, 2011 9:16 am Post subject: osgCompute / osgCuda Postfilter |
|
|
Hi,
First of all thanks for the great work you've benn doing so far. I'm trying to build up a post processing setup using osgCompute but I have some problems.
This is how my scene graph looks like:
Computation (Render before children)
|-> Camera (Relative, attached MRT textures and using FBO)
|........|--> Scene
|-> Camera (Absolute, Orthogonal 2D Camera)
.........|--> Quad Geometry using Target Texture
I'm generating the MRTs and Target textures as follows:
Code:
|
char MrtName[80];
for(int i = 0; i < 4; i++)
{
MRTTexture[i] = new osgCuda::Texture2D;
MRTTexture[i]->setTextureSize(textureWidth, textureHeight);
MRTTexture[i]->setFilter(osg::Texture2D::MIN_FILTER,osg::Texture2D::LINEAR);
MRTTexture[i]->setFilter(osg::Texture2D::MAG_FILTER,osg::Texture2D::LINEAR);
MRTTexture[i]->setInternalFormat(GL_RGBA32F_ARB);
MRTTexture[i]->setSourceFormat(GL_RGBA);
MRTTexture[i]->setSourceType(GL_FLOAT);
sprintf(MrtName, "MRT%d", i);
MRTTexture[i]->addIdentifier(MrtName);
}
//Target
*Target = new osgCuda::Texture2D;
(*Target)->setInternalFormat( GL_RGBA32F_ARB );
(*Target)->setSourceFormat( GL_RGBA );
(*Target)->setSourceType( GL_FLOAT );
(*Target)->setTextureWidth( textureWidth );
(*Target)->setTextureHeight( textureHeight );
(*Target)->setFilter(osg::Texture2D::MIN_FILTER, osg::Texture2D::NEAREST);
(*Target)->setFilter(osg::Texture2D::MAG_FILTER, osg::Texture2D::NEAREST);
(*Target)->addIdentifier( "TARGET" );
|
And this is how I call my kernel:
Code:
|
virtual void launch()
{
if( isClear() )
return;
swap( _blocks,
_threads,
Target->map(osgCompute::MAP_DEVICE_TARGET),
MRT0->map(osgCompute::MAP_DEVICE_SOURCE),
Target->getPitch(),
MRT0->getPitch(),
Target->getDimension(0),
Target->getDimension(1) );
}
|
And this is the kernel:
Code:
|
__global__
void swapKernel( float4* dst, float4* src, unsigned int dstPitch, unsigned int srcPitch, unsigned int imageWidth, unsigned int imageHeight )
{
// compute thread dimension
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
if( x < imageWidth && y < imageHeight )
{
float4* dstPixel = (float4*)(((char*) dst) + dstPitch * y ) + x;
float4* srcPixel = (float4*)(((char*) src) + srcPitch * y ) + x;
// swap channels
*dstPixel = *srcPixel;
//*dstPixel = make_float4( ((float) x / (float) imageWidth) ,((float) y / (float) imageHeight), 0, 1 );
}
}
|
I'm sure MRT's are setup correctly since when I attach any of the MRT's to the Quad I can see the scene.
I'm also sure that kernel is working properly since when I write texture coordinates to the target I can see them.
However when I try to copy the MRT to the Target I see a black screen.
I couldn't find whats wrong with this setup. Any help would be appreciated.
Thanks in advance!
Cheers,
Serkan[/code] |
|
Back to top |
|
 |
J.P. Delport Guest
|
Posted: Tue Feb 01, 2011 1:25 pm Post subject: osgCompute / osgCuda Postfilter |
|
|
Hi,
I've recently been fiddling/debugging with a similar setup, nl RTT to
osgCuda. From what I've seen in the code (r314) the copy step is missing
for osgCompute::MAP_DEVICE_SOURCE. Allocation of a copy happens, but if
the texture does not have an osg::Image, no sync is taking place.
Sending my test app to the osgCuda devs is on my todo list that does not
seem to want to get shorter. I was also not sure if I have to attach
some kind of callback.
cheers
jp
On 01/02/11 11:16, Serkan Ergun wrote:
Quote:
|
Hi,
First of all thanks for the great work you've benn doing so far. I'm trying to build up a post processing setup using osgCompute but I have some problems.
This is how my scene graph looks like:
Computation (Render before children)
|-> Camera (Relative, attached MRT textures and using FBO)
|........|--> Scene
|-> Camera (Absolute, Orthogonal 2D Camera)
.........|--> Quad Geometry using Target Texture
I'm generating the MRTs and Target textures as follows:
Code:
char MrtName[80];
for(int i = 0; i< 4; i++)
{
MRTTexture[i] = new osgCuda::Texture2D;
MRTTexture[i]->setTextureSize(textureWidth, textureHeight);
MRTTexture[i]->setFilter(osg::Texture2D::MIN_FILTER,osg::Texture2D::LINEAR);
MRTTexture[i]->setFilter(osg::Texture2D::MAG_FILTER,osg::Texture2D::LINEAR);
MRTTexture[i]->setInternalFormat(GL_RGBA32F_ARB);
MRTTexture[i]->setSourceFormat(GL_RGBA);
MRTTexture[i]->setSourceType(GL_FLOAT);
sprintf(MrtName, "MRT%d", i);
MRTTexture[i]->addIdentifier(MrtName);
}
//Target
*Target = new osgCuda::Texture2D;
(*Target)->setInternalFormat( GL_RGBA32F_ARB );
(*Target)->setSourceFormat( GL_RGBA );
(*Target)->setSourceType( GL_FLOAT );
(*Target)->setTextureWidth( textureWidth );
(*Target)->setTextureHeight( textureHeight );
(*Target)->setFilter(osg::Texture2D::MIN_FILTER, osg::Texture2D::NEAREST);
(*Target)->setFilter(osg::Texture2D::MAG_FILTER, osg::Texture2D::NEAREST);
(*Target)->addIdentifier( "TARGET" );
And this is how I call my kernel:
Code:
virtual void launch()
{
if( isClear() )
return;
swap( _blocks,
_threads,
Target->map(osgCompute::MAP_DEVICE_TARGET),
MRT0->map(osgCompute::MAP_DEVICE_SOURCE),
Target->getPitch(),
MRT0->getPitch(),
Target->getDimension(0),
Target->getDimension(1) );
}
And this is the kernel:
Code:
__global__
void swapKernel( float4* dst, float4* src, unsigned int dstPitch, unsigned int srcPitch, unsigned int imageWidth, unsigned int imageHeight )
{
// compute thread dimension
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
if( x< imageWidth&& y< imageHeight )
{
float4* dstPixel = (float4*)(((char*) dst) + dstPitch * y ) + x;
float4* srcPixel = (float4*)(((char*) src) + srcPitch * y ) + x;
// swap channels
*dstPixel = *srcPixel;
//*dstPixel = make_float4( ((float) x / (float) imageWidth) ,((float) y / (float) imageHeight), 0, 1 );
}
}
I'm sure MRT's are setup correctly since when I attach any of the MRT's to the Quad I can see the scene.
I'm also sure that kernel is working properly since when I write texture coordinates to the target I can see them.
However when I try to copy the MRT to the Target I see a black screen.
I couldn't find whats wrong with this setup. Any help would be appreciated.
Thanks in advance!
Cheers,
Serkan[/code]
------------------
Read this topic online here:
http://forum.openscenegraph.org/viewtopic.php?p=36216#36216
|
--
This message is subject to the CSIR's copyright terms and conditions, e-mail legal notice, and implemented Open Document Format (ODF) standard.
The full disclaimer details can be found at http://www.csir.co.za/disclaimer.html.
This message has been scanned for viruses and dangerous content by MailScanner,
and is believed to be clean. MailScanner thanks Transtec Computers for their support.
------------------
Post generated by Mail2Forum |
|
Back to top |
|
 |
serkan.ergun Newbie
Joined: 01 Feb 2011 Posts: 3
|
Posted: Wed Feb 02, 2011 7:52 am Post subject: Re: osgCompute / osgCuda Postfilter |
|
|
J.P. Delport wrote:
|
Hi,
I've recently been fiddling/debugging with a similar setup, nl RTT to
osgCuda. From what I've seen in the code (r314) the copy step is missing
for osgCompute::MAP_DEVICE_SOURCE. Allocation of a copy happens, but if
the texture does not have an osg::Image, no sync is taking place.
Sending my test app to the osgCuda devs is on my todo list that does not
seem to want to get shorter. I was also not sure if I have to attach
some kind of callback.
cheers
jp
|
Hi,
Is there any fix or workaround for this that you've found?
cheers
Serkan |
|
Back to top |
|
 |
J.P. Delport Guest
|
Posted: Wed Feb 02, 2011 8:25 am Post subject: osgCompute / osgCuda Postfilter |
|
|
Hi,
On 02/02/11 09:52, Serkan Ergun wrote:
Quote:
|
J.P. Delport wrote:
Quote:
|
Hi,
I've recently been fiddling/debugging with a similar setup, nl RTT to
osgCuda. From what I've seen in the code (r314) the copy step is missing
for osgCompute::MAP_DEVICE_SOURCE. Allocation of a copy happens, but if
the texture does not have an osg::Image, no sync is taking place.
Sending my test app to the osgCuda devs is on my todo list that does not
seem to want to get shorter. I was also not sure if I have to attach
some kind of callback.
cheers
jp
|
Hi,
Is there any fix or workaround for this that you've found?
|
No, I haven't had time to look at this further. I expect that a
cudaMemCpy from the RTT texture to the allocated osgCuda copy will fix it.
jp
Quote:
|
cheers
Serkan
------------------
Read this topic online here:
http://forum.openscenegraph.org/viewtopic.php?p=36253#36253
|
--
This message is subject to the CSIR's copyright terms and conditions, e-mail legal notice, and implemented Open Document Format (ODF) standard.
The full disclaimer details can be found at http://www.csir.co.za/disclaimer.html.
This message has been scanned for viruses and dangerous content by MailScanner,
and is believed to be clean. MailScanner thanks Transtec Computers for their support.
------------------
Post generated by Mail2Forum |
|
Back to top |
|
 |
jens.svt User
Joined: 16 Mar 2009 Posts: 30
|
Posted: Fri Feb 04, 2011 4:05 pm Post subject: Solution for processing GL render targets |
|
|
Hi serkan, hi J.P.,
Here is a link to a simple example solving your problem:
http://www.cg.informatik.uni-siegen.de/data/Downloads/svt/osgComputeDemoApp.zip
We render a rotating cow into a texture then run a cuda-kernel and display the result in the same scene.
You need to set the usage flags of your render targets to GL_TARGET, e.g. :
Target0->setUsage( osgCompute::GL_TARGET_COMPUTE_SOURCE ).
With this flag you tell the texture that it is used as a FBO.
Unfortunately, OpenSceneGraph does not call any function of osg::Texture when it renders into a texture.
In other words something like texture->applyAsFBO() is missing.
So there is no way to get notified when this happens.
The only way to deal with this is to copy the texture memory each time a mapping function is called (GPU->GPU which is still fast).
However, for this the user needs to setup the usage flag.
Additionally, with older drivers/CUDA versions (we are not sure about this) you have to call MRTTexture->unmap() after your kernel has finished.
If you do not call unmap(), the MRTTexture is still mapped in the CUDA context and OpenGL cannot render into it successfully.
Hopefully this helps,
Jens |
|
Back to top |
|
 |
J.P. Delport Guest
|
Posted: Mon Feb 07, 2011 8:32 am Post subject: osgCompute / osgCuda Postfilter |
|
|
Hi Jens,
On 04/02/11 18:05, Jens Orthmann wrote:
Quote:
|
Hi serkan, hi J.P.,
Here is a link to a simple example solving your problem:
http://www.cg.informatik.uni-siegen.de/data/Downloads/svt/osgComputeDemoApp.zip
We render a rotating cow into a texture then run a cuda-kernel and display the result in the same scene.
You need to set the usage flags of your render targets to GL_TARGET, e.g. :
Target0->setUsage( osgCompute::GL_TARGET_COMPUTE_SOURCE ).
With this flag you tell the texture that it is used as a FBO.
Unfortunately, OpenSceneGraph does not call any function of osg::Texture when it renders into a texture.
In other words something like texture->applyAsFBO() is missing.
So there is no way to get notified when this happens.
The only way to deal with this is to copy the texture memory each time a mapping function is called (GPU->GPU which is still fast).
However, for this the user needs to setup the usage flag.
|
I added the flag and my test app is now working, thanks.
regards
jp
Quote:
|
Additionally, with older drivers/CUDA versions (we are not sure about this) you have to call MRTTexture->unmap() after your kernel has finished.
If you do not call unmap(), the MRTTexture is still mapped in the CUDA context and OpenGL cannot render into it successfully.
Hopefully this helps,
Jens
------------------
Read this topic online here:
http://forum.openscenegraph.org/viewtopic.php?p=36373#36373
|
--
This message is subject to the CSIR's copyright terms and conditions, e-mail legal notice, and implemented Open Document Format (ODF) standard.
The full disclaimer details can be found at http://www.csir.co.za/disclaimer.html.
This message has been scanned for viruses and dangerous content by MailScanner,
and is believed to be clean. MailScanner thanks Transtec Computers for their support.
------------------
Post generated by Mail2Forum |
|
Back to top |
|
 |
serkan.ergun Newbie
Joined: 01 Feb 2011 Posts: 3
|
Posted: Mon Feb 07, 2011 8:59 am Post subject: |
|
|
Hi,
Thank you very much! This solved my problem.
Cheers,
Serkan |
|
Back to top |
|
 |
sholmes User

Joined: 05 Sep 2011 Posts: 62
|
Posted: Tue May 01, 2012 1:04 pm Post subject: Re: Solution for processing GL render targets |
|
|
I am trying to mod this sample to use the following texture type... Does osgCompute support this?
textureRect0->setSourceType(GL_FLOAT);
textureRect0->setInternalFormat(GL_RGBA16F_ARB);
textureRect0->setSourceFormat(GL_RGBA);
jens.svt wrote:
|
Hi serkan, hi J.P.,
Here is a link to a simple example solving your problem:
http://www.cg.informatik.uni-siegen.de/data/Downloads/svt/osgComputeDemoApp.zip
We render a rotating cow into a texture then run a cuda-kernel and display the result in the same scene.
You need to set the usage flags of your render targets to GL_TARGET, e.g. :
Target0->setUsage( osgCompute::GL_TARGET_COMPUTE_SOURCE ).
With this flag you tell the texture that it is used as a FBO.
Unfortunately, OpenSceneGraph does not call any function of osg::Texture when it renders into a texture.
In other words something like texture->applyAsFBO() is missing.
So there is no way to get notified when this happens.
The only way to deal with this is to copy the texture memory each time a mapping function is called (GPU->GPU which is still fast).
However, for this the user needs to setup the usage flag.
Additionally, with older drivers/CUDA versions (we are not sure about this) you have to call MRTTexture->unmap() after your kernel has finished.
If you do not call unmap(), the MRTTexture is still mapped in the CUDA context and OpenGL cannot render into it successfully.
Hopefully this helps,
Jens
|
|
|
Back to top |
|
 |
jens.svt User
Joined: 16 Mar 2009 Posts: 30
|
Posted: Wed May 02, 2012 12:58 pm Post subject: |
|
|
Hi sholmes,
I use GL_RGBA32F in my projects.
However, CUDA had problems to map
GL_RGBA16F in earlier versions.
Maybe you should try to run this internal GL format
with our RTT-example first.
Cheers,
Jens |
|
Back to top |
|
 |
sholmes User

Joined: 05 Sep 2011 Posts: 62
|
Posted: Wed May 02, 2012 5:03 pm Post subject: |
|
|
I was successfull in using the GLRGBA32F as well as the GL_LUMINANCE_ALPHA32F. Thanks for the feedback. osgCompute is a very interesting project... Would like to see across the board support for multithreaded mode though... Is that in the works?
jens.svt wrote:
|
Hi sholmes,
I use GL_RGBA32F in my projects.
However, CUDA had problems to map
GL_RGBA16F in earlier versions.
Maybe you should try to run this internal GL format
with our RTT-example first.
Cheers,
Jens
|
|
|
Back to top |
|
 |
|
|
You cannot post new topics in this forum You cannot reply to topics in this forum You cannot edit your posts in this forum You cannot delete your posts in this forum You cannot vote in polls in this forum You cannot attach files in this forum You cannot download files in this forum
|
|
Powered by phpBB © 2001, 2005 phpBB Group
|