OpenSceneGraph Forum Forum Index OpenSceneGraph Forum
Official forum which mirrors the existent OSG mailing lists. Messages posted here are forwarded to the mailing list and vice versa.
 
   FAQFAQ    SearchSearch    MemberlistMemberlist    RulesRules    UsergroupsUsergroups    RegisterRegister 
 Mail2Forum SettingsMail2Forum Settings  ProfileProfile   Log in to check your private messagesLog in to check your private messages   Log inLog in 
   AlbumAlbum  OpenSceneGraph IRC ChatOpenSceneGraph IRC Chat   SmartFeedSmartFeed 

osgCompute / osgCuda Postfilter


 
Post new topic   Reply to topic    OpenSceneGraph Forum Forum Index -> osgCompute [osgCompute]
View previous topic :: View next topic  
Author Message
serkan.ergun
Newbie


Joined: 01 Feb 2011
Posts: 3

PostPosted: Tue Feb 01, 2011 9:16 am    Post subject:
osgCompute / osgCuda Postfilter
Reply with 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]
Back to top
View user's profile Send private message
J.P. Delport
Guest





PostPosted: Tue Feb 01, 2011 1:25 pm    Post subject:
osgCompute / osgCuda Postfilter
Reply with 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

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

PostPosted: Wed Feb 02, 2011 7:52 am    Post subject:
Re: osgCompute / osgCuda Postfilter
Reply with quote

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
View user's profile Send private message
J.P. Delport
Guest





PostPosted: Wed Feb 02, 2011 8:25 am    Post subject:
osgCompute / osgCuda Postfilter
Reply with quote

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

PostPosted: Fri Feb 04, 2011 4:05 pm    Post subject:
Solution for processing GL render targets
Reply with 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.

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
View user's profile Send private message
J.P. Delport
Guest





PostPosted: Mon Feb 07, 2011 8:32 am    Post subject:
osgCompute / osgCuda Postfilter
Reply with quote

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

PostPosted: Mon Feb 07, 2011 8:59 am    Post subject:
Reply with quote

Hi,

Thank you very much! This solved my problem.

Cheers,
Serkan
Back to top
View user's profile Send private message
sholmes
User


Joined: 05 Sep 2011
Posts: 62

PostPosted: Tue May 01, 2012 1:04 pm    Post subject:
Re: Solution for processing GL render targets
Reply with quote

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
View user's profile Send private message
jens.svt
User


Joined: 16 Mar 2009
Posts: 30

PostPosted: Wed May 02, 2012 12:58 pm    Post subject:
Reply with quote

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
View user's profile Send private message
sholmes
User


Joined: 05 Sep 2011
Posts: 62

PostPosted: Wed May 02, 2012 5:03 pm    Post subject:
Reply with quote

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
View user's profile Send private message
Display posts from previous:   
Post new topic   Reply to topic    OpenSceneGraph Forum Forum Index -> osgCompute [osgCompute] All times are GMT
Page 1 of 1

 
Jump to:  
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

Similar Topics
Topic Author Forum Replies Posted
No new posts How to postprocess the data in depth ... eastwood osgCompute [osgCompute] 3 Mon Aug 25, 2014 4:56 pm View latest post
No new posts is osgCompute still active? sholmes General [forum] 1 Tue Jul 08, 2014 1:49 pm View latest post
No new posts osgCompute Questions [or... osgComput... sholmes osgCompute [osgCompute] 0 Wed Apr 09, 2014 2:10 pm View latest post
No new posts osgCompute Questions Conan General 0 Mon Apr 07, 2014 6:52 pm View latest post
No new posts osgCompute example SMesserschmidt General 1 Wed Dec 04, 2013 9:44 am View latest post


Board Security Anti Bot Question MOD - phpBB MOD against Spam Bots
Powered by phpBB © 2001, 2005 phpBB Group
Protected by Anti-Spam ACP