Newton on GPU (CUDA, OpenCL)

[Julio Jerez]

Here is a small previews of the new pawaaaaaaaarrrrrrr
http://www.newtondynamics.com/downloads ... Solver.wmv

This is not optimized and it only implements the solver in the GPU.
Broad Phase, narrow phase collision, island assembly and body integration are still implemented in the CPU.
What is worse I the and the beginning of the solver the entire scene is serialized to the GPU on entry on serialize to the CPU on exit.
Even with all these ineficioencies you can see that the CPU single core 2.66 core2 pentium, the profile is off the crart, while the GPU ge 280 is about 30 fps.

As the other system have their own proxy representation in GPU memory all the memory copy will go away and the how system should be way, way faster.
The reason I stared this in reverse order (solver firt) in because the Newton solver is the more critical part of the engine, and if this part do not work, then all the rest is useless, This is an experience I had when trying to implement the Open MP version and the impossibility of efficient thread management.
Any way I will continue adding all of the major system, and then I will move to the simpler and trial stuff like, cloth, soft bodies, and SPH fluids.


This works on all CUDA ready GForce.
In fact this can even be implemented in a pixel Shader but the is some thing I will not dare doing.
The Cuda version will be a temperoraty GPU implemenation until the more generic OpenCL language is ready, and then will be abailable on all GPU, and Muticores CPU.
I can not way for the Larrabee which I think will smoke big the Gforce to teh dust.

[VeT]

wow... how many blocks are there?

//Any way I will continue adding all of the major system, and then I will move to the simpler and trial stuff like, cloth, soft bodies, and SPH fluids.
yeah!

[agi_shi]

Can't wait for a new beta ...

[Julio Jerez]

Solver optimization really makes a difference with CUDA.
In the inner loop of the more frequent called function I have the dot product calculation

Code: Select all

rowAccel.m_jacobianData[thread] = jacobianMemory[first].m_JMinv.m_jacobianData[thread] * internalForce.m_jacobianData[thread];
__syncthreads();   
for (int n = 8; n; n >>= 1) {      
        if (thread < n) {   
             rowAccel.m_jacobianData[thread] += rowAccel.m_jacobianData[thread + n];
       }
      __syncthreads();         
}

This is a dot product and would be the correct way of doing on any hardware, However the Gforce is not a very well designed CPU and the programmer have o take care of lot of short coming in the hardware
So by doing a circular sum which will use at least twice as many operation the result in quite significant in performance for just one instruction in a much large function
Changed the code to this:

Code: Select all

for (int n = 8; n; n >>= 1) {      
     rowAccel.m_jacobianData[thread] += rowAccel.m_jacobianData[(thread + n) & 15];
}

}

Here is a video of the different of just that simple line
http://www.newtondynamics.com/downloads ... tepOne.wmv

I still have figure out how to get the occupancy higher, according to the profiler all of my kernels are only using 25% of the capacity of the multiprocessor.
Thsi is good because it means teher is a lot of room for optimization.
If only nvidia people answer my questions, I posted about tow weeks ago, and I only get view but no one answer.
http://forums.nvidia.com/index.php?showtopic=88445

Maybe the also do not want us making waves or embarras physics CUDA apps, like the Collada leaders told me.

anyway I still have a lot to do implemention the other faces of the engine with will eliminate the memory big copies, then will will have real sustancial gains.
the Hardware optimizatiosn I leave for later.

[Sury]

I can not way for the Larrabee which I think will smoke big the Gforce to teh dust.

At NVISION08 they call Larrabee "marketing puff" and told the press that the Larrabee architecture was "like a GPU from 2006" . I like the idea behind it. Seems to be much more flexible. Probably less than a year to see it on the market.

[Julio Jerez]

NVISION? by the N in the front, i am guessing this is another Nvidia acquired product, promoted by nvidia propaganda machine.
I would not to put too much credibility on nvidia spokesman about the competition. Nvidia has many that are not very trustworthy.

any Pentium is by far much, much superior than the more advance GPU of today.
when you look under the hook a GPU is quite a mediocre and limited piece of hardware.
Larrabee will smoke them all big, big time.

Saying that an nvidia GPU is better than a pemtion CPU or a AMD general CPU, is the equivalent to say a Nasca Car is better than a Rolls Royce, just because the race car can do faster laps on a circular track.
I am sure some Nasca fanatic will believe it, but try to sell that to any normal buyer.
the people claiming nVidia superiority are the people associated with nVidia either because they work for them or because nvidia bought them out,
and the unsolicited cheerleader in forums and internet technology news outlets that are heavy funded by Nvidia ads and other kind promotional gifts.

[Sury]

Here is the article.
http://www.pcpro.co.uk/news/220947/nvis ... -2006.html

If I was nVidia, I would affraid of CPU mosters like Intel. If i get it right, this thing will be a small general purpose multicore x86 thingy with more than 32 cores running at over 1 GHz, communicating over a 1024 bit wide buss, each core is 4 way multithreaded,architecture has a pipeline derived from the dual-issue Intel Pentium and stuff. As they say, real-time raytracing is not a dream any more and that yelds to very cool renderings and effects. Pixar and DreamWorks are interested

[Sury]

I don't know...probably all hype about this or that hardware is going to end up in a higher than the end user level. If you wanna play games , just plug your old PC to a fast broadband connection and have instant fun.
See what's comming along
http://www.onlive.com/index.html

[Stucuk]

Looks interesting, tho there is only mention of bringing it to the US. I don't get why US companies always just release in the US, then 50 years later release in other countries. Any company which isn't US based would release world wide and think of an international scale. Unsupprisingly there is a larger amount of people who don't live in the US, meaning world wide distribution would mean far more money for them than just US only.

[Julio Jerez]

It look like nvidia is testing and OpenCL driver
http://www.nvidia.com/object/io_1240224603372.html

has anyone try it.
I registered long time ago and I have got anything

I have a working CUDA Version on teh engine that I can refine and optimize,
But given all the toumoir with OpenCL, Larrabee, CUDA, and ATI,
It is kind of hard to make such huge effort and then have all of it reder useless just because Nvdia decide to change CUDA unilaterally.
I like this multiplatform aspect of OpenCL.

[djbe]

Hey julio,

I myself have been quite interested in OpenCL, would you happen to have any good documentation about it (something like simple examples to get started) and any info on how to use it now? (Like what software/libraries do I need, where can I get them, etc...).
Cya!

[Dave Gravel]

They seen to have do some OpenCL demos already.
http://developer.amd.com/SAMPLES/STREAM ... fault.aspx

[Julio Jerez]

I hav eseen those videos in you tube.
They say it is OpenCl, but then when you go to ATI, they do no have any public sopport for OpenCL.
If they do they are keeping very secrect, which I think is misinformation.

All I see at the ATI site is Stream SDK stuff, which is another proprietary language.

[Dave Gravel]

Ok yes possible me too I have seen it from youtube first, I have find it by hassard.

[vartaxe]

after reading this http://www.fudzilla.com/content/view/14271/1/ i have to say opencl is the way to go