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 

Persistant texture modification using osgCompute / CUDA


 
Post new topic   Reply to topic    OpenSceneGraph Forum Forum Index -> osgCompute [osgCompute]
View previous topic :: View next topic  
Author Message
Asmar Arsala
Guest





PostPosted: Wed Nov 04, 2009 12:05 pm    Post subject:
Persistant texture modification using osgCompute / CUDA
Reply with quote

Dear all,

I’m interested to see whether I can use osgCompute for my project’s goal, which is real-time height-map deformation system. For starters, I’m trying to let an osgCompute node modify a 2D OSG float texture (the height-map) each frame, using a CUDA kernel.
My problem now is that I can’t get osgCompute to make persistent changes to the source texture, it seems that either the CUDA kernel changes the texture only once or that it performs the same changes each frame to the original texture data. Any advice or hints on how to set up my buffers, textures and arrays in osgCompute to achieve persistent changes would be greatly appreciated.
I’m using osgCompute 0.4 (with CUDA 2.3, OSG 2.8.2, VS 9.0 / WinXP) and I started off with modifying the provided osgTexDemo example:

With regard to the original source code I changed the SrceBuffer and TempBuffer to single dimension (linear) osg::Cuda::Arrays with the intent to alternatingly read from, and write to, these respectively.
The extended (overloaded) osgCompute::Module::launch() alternatingly maps the SrceBuffer and TempBuffer as the source and target for the primary Cuda kernel respectively.
The setImage() call in the else is a hack to ensure that osgCompute cannot overwrite the buffer with the original image again when mapping the buffer to the device. This code so far has been tailored to simply try get it to work based on the original example; I can image this not being to best approach however.

Sincerely,
Asmar B. Arsala

(for code snippets see attachment) Express yourself instantly with MSN Messenger! MSN Messenger

------------------
Post generated by Mail2Forum
Back to top
jens.svt
User


Joined: 16 Mar 2009
Posts: 30

PostPosted: Thu Nov 05, 2009 3:12 pm    Post subject:
Reply with quote

Hi Asmar,

I have taken a look at your code fragment and i have found out that you write into osgCuda::Arrays. This might cause your problem because the "osgCuda::Array" class is an abstraction for "cudaArray" objects which are read-only. Instead you should use "osgCuda::Buffer" objects into which kernels are allowed to write.

In the current version of osgCuda you get an error message when mapping the array as a TARGET for the DEVICE.

Cheers,
Jens
Back to top
View user's profile Send private message
Asmar Arsala
Guest





PostPosted: Mon Nov 16, 2009 12:47 pm    Post subject:
Persistant texture modification using osgCompute / CUDA
Reply with quote

Hey Jens,
<?xml:namespace prefix = o ns = "urn:schemas-microsoft-com:office:office" />
Thanks for you quick reply. As I have been away from the project for a while due to illness I hadn’t been able to really act upon your comments until now.

It had also been a while since I worked with Cuda and I had forgotten that Cuda actually makes a crucial distinction between memory allocated through cudaMalloc() versus memory allocated by cudaMallocArray(). In my previous attempts I made the erroneous assumption that a one-dimensional array through osgCuda::Array would lead to cuda-writable memory.

After perusing the osgCuda and osgCompute code I now gather that osgCuda::Array allocates with cudaMallocArray(), osgCuda::Buffer allocates with cudaMalloc() and osgCuda::Texture allocates via OpenGL Buffer Objects. Which would indeed suggest that osgCuda::Arrays are read only, osgCuda::Buffers are read-write, and osgCuda::Textures are read-write (depending on the current OpenGL mapping and binding states). I should point out here that my background is DirectX and that I am not yet all that wise in the ways of OpenGL in general.

I had originally tried using osgCuda::Buffers, but then realised that the cudaBindTextureToArray() was failing (for obvious reasons). I then changed to one dimensional osgCuda::Arrays under the aforementioned assumption. In retrospect after having the source data be linear, there wasn’t really any need to map to texture memory, so I could have just as easily stayed with the osgCuda::Buffers and passed both source and destination data pointers to the kernel. This indeed, works fine, however the whole setup has now kind of strayed from the initial intent.

The idea is to be able to persistently update textures containing height related data on numerous LOD levels. Hence I want to avoid memory copies as much as possible; especially from host to device, and vice verse.
Now I also gather that the both osgCuda::Arrays and osgCuda::Buffers, due to their nature should be used as very temporary buffers; for cuda computations and not as persistent (texture-related) memory. Hence osgCuda::Textures are in fact the logical choice here. OpenGL can take care of the (device) memory management, Cuda can have read and write access, and changes can remain persistent. I imagine one of the most efficient ways to update a texture would then be to first map to a cudaArray for reading and then writing the result back to the osgCuda::Texture PBO.

To test this out I have now changed all osgCuda::Buffers/Arrays to osgCuda::Textures, however this almost always throws a number of exceptions; sometimes when I debug it (without changing the code) is does not throw these exceptions. When that happens it will actually run perfectly. Otherwise it just doesn’t work. The exceptions thrown are:

osgCuda::Context::registerBufferObject() for Context "0": something goes wrong within cudaGLRegisterBufferObject().
unspecified launch failure.
osgCuda::Texture2D::allocPBO() for Buffer "TempBuffer": Could not generate buffer object (glGenBuffers()) for context "0".
osgCuda::Texture::mapStream() for texture "TempBuffer": error during cudaGLMapBufferObject() for context "0".
unspecified launch failure.
osgCuda::Context::registerBufferObject() for Context "0": something goes wrong within cudaGLRegisterBufferObject().
unspecified launch failure.
osgCuda::Texture2D::allocPBO() for Buffer "DestBuffer": Could not generate buffer object (glGenBuffers()) for context "0".
osgCuda::Texture::mapStream() for texture "DestBuffer": error during cudaGLMapBufferObject() for context "0".
unspecified launch failure.

That is, something goes wrong in Context::mallocBufferObject().

Now I’m starting to suspect that the way osgCompute is set up might not be ideal for what I intend to do here, but I would also not be the least bit surprised if I am totally misusing something again.
As before, any suggestions are greatly appreciated.

Sincerely,
Asmar B. Arsala Express yourself instantly with MSN Messenger! MSN Messenger

------------------
Post generated by Mail2Forum
Back to top
Asmar Arsala
Guest





PostPosted: Fri Nov 20, 2009 4:24 pm    Post subject:
Persistant texture modification using osgCompute / CUDA
Reply with quote

Hey Jens,

You can pretty much ignore the previous mail I sent; I ended up upgrading to the SVN version and I adapted the new TexDemo example to read from and write to a single osgCudaTexture, which seems to work fine. Thanks.

Sincerely,
Asmar

P.S.

By the by, in the gaussKernel() of the osgTextDemo, I am assuming your intent is to normalize the src vector when you divide it by 423.0f, the sum of the current weights is however 577.0f. Smile
Express yourself instantly with MSN Messenger! MSN Messenger

------------------
Post generated by Mail2Forum
Back to top
jens.svt
User


Joined: 16 Mar 2009
Posts: 30

PostPosted: Fri Nov 20, 2009 5:27 pm    Post subject:
Reply with quote

Hi Asmar,

that is great because we were not able to reproduce the error. Shocked
Thanks a lot for the wrong normalization factor.

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 obj plugin does not support diffuse ... Ralf Habacker Submission 0 Tue Nov 14, 2017 4:18 pm View latest post
No new posts Get RGB Values from Texture dieslower81 General 6 Tue Nov 07, 2017 3:56 pm View latest post
No new posts Problem With Osg Texture Colors! Bobykhani General [forum] 1 Tue Oct 10, 2017 8:20 am View latest post
No new posts obj plugin does not support diffuse a... Ralf Habacker General 5 Tue Sep 19, 2017 6:21 am View latest post
No new posts How to texture map Box/Sphere/Cone/Ca... pleopard General 2 Mon Aug 07, 2017 7:14 pm 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