OpenSceneGraph Forum

[serkan.ergun]

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]

[J.P. Delport]

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

[serkan.ergun]

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

[J.P. Delport]

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

[jens.svt]

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

[J.P. Delport]

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

[serkan.ergun]

Hi,

Thank you very much! This solved my problem.

Cheers,
Serkan

[sholmes]

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

[jens.svt]

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

[sholmes]

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