FDS and CUDA

[Hendrik]

I did some modifications to the source code and added a module written
in C/CUDA to FDS. I got some fine results using my NVIDIA GTX 280 card
as shown at

www.qfds.de.

Although I know that you probably can not add this to your project
(sice the modifications are not written in FORTRAN) I wanted to post
it anyway, just in case someboby is interested.

Regards,
Hendrik

[Christian Rogsch]

Hendrik,

I think these results are very important for speed-up settings in FDS.
I think I is possible to write code in that way that the writing
routines could be run in parallel to the calculation routines, thus we
have no speed-limitation based on writing to disk.

I think we should discuss this possibility next week at the meeting in
Berlin.

Regards,
Christian

[Glenn Forney]

I am interested in this. Could you explain in a little more detail what
changes you made? Are you replacing an FDS subroutine with a CUDA
equivalent?

--
Glenn Forney
National Institute of Standards and Technology
100 Bureau Drive, Stop 8663
Gaithersburg MD 20899-8663

Telephone: (301) 975 2313
FAX: (301) 975 4052

Pre-decisional and sensitive information. Not for attribution, distribution, or reproduction.

[Dave McGill]

Hello All,

We have a brand new NVIDIA Tesla S1070 (four cards, and the front node
runs Centos 5.2 64bit with CUDA sdk). See, http://www.nvidia.com/object/product_tesla_s1070_us.html
. If access to this would be of any value in this development work,
let me know and I'll arrange to get accounts set up.

Regards

Dave McGill
Seneca College

[Hendrik]

Hi Christian,

I just did'nt had the time to transfer the CPU-intensive parts of the
DUMP modul to CUDA yet (i.e. SUBROUTINE UPDATE_HHR) . We'll see what's
left of it, when I've done this.

See you next week...
Hendrik

[Hendrik]

Hello Glenn,

yes, indeed. I replaced every loop in FORTRAN by an equivalent CUDA-
kernel in C like this one:

__global__ void mass_kernel_cd4(int nDimI, int nDimJ, int nDimK, REAL
lfRHOMAX, REAL *WORK1 /*RHODELTA*/, REAL *RHOP)
{
FDS_IJK_0
RHOP(I,J,K) = FMIN(lfRHOMAX, RHOP(I,J,K) + RHODELTA(I,J,K));
}

In addition to that I did some changes in the data structures (i.e.
you're going to need more temporary arrays, if you want to compute the
heat conduction in more than one wall cell at a time...).

So unfortunately it is not this easy to implement as Christian showed
in his OpenMP sample...

[Hendrik]

Hi,

thank you for the offer. Up until now I can only use one GPU at a
time, so I guess one C1060 would be sufficient so far. But I would
like to come to your offer once the development is more processed.

Regards
Hendrik

On Feb 26, 7:30 pm, Dave McGill <david.mcg...@senecac.on.ca> wrote:
> Hello All,
>
> We have a brand new NVIDIA Tesla S1070 (four cards, and the front node

> runs Centos 5.2 64bit with CUDA sdk). See,http://www.nvidia.com/object/product_tesla_s1070_us.html

[Dave McGill]

Hi Hendrik,

That's fine. Just let me know when you want an account.

Regards

Dave

[Kevin]

This is very interesting. Just for your information, we are planning
to implement the OpenMP calls that Christian has suggested. This seems
like the easiest way to exploit multi-processor, multi-core chips. But
we will continue to explore other ways to parallelize, like the one
you have suggested. A question for you -- would the OpenMP calls make
it more difficult to work with the GPU? I know very little about this
subject.

Thanks

[Hendrik]

For now I am working in C and I don't know what the FORTRAN version of
CUDA will look like. But in my opinion it should be possible - once
CUDA/FORTRAN is released - to add new "GPU-lines" around the loops
similar to OpenMP. You only have to write a little tool which replaces
the "GPU-Lines" by adequate calls to CUDA kernels before compiling the
source code. Maybe even the OpenMP calls could be used for this.

I have a question too -- in the user manual is written you did some
tests using single precision variables. Do you have any material
concerning this matter left? Since the GPU's are about 8 times faster
in single precision mode I would like to see if its possible to carry
out parts of the computations using float numbers (i.e. pressure
solver, heat conduction in solids, radiation... ).

Thanks

[Kevin]

I once tried to run a very early version of the FDS solver with single
precision (4 byte reals or what used to be called REAL*4 in Fortran
77). The problem was the pressure solver. We use a direct (non-
iterative) solver for the Poisson equation, meaning that the accuracy
of the discrete solution is roughly on the order of 10^-12 in the
residuals. For very long calculations, this error builds up, but to
levels that are still very small (like 10^-7, for example). In single
precision, the error build-up can actually cause problems, at least it
did when I tried it many years ago. I haven't tried since, mainly
because we're all moving towards 64 bit (8 byte) arithmetic on 64 bit
OS platforms as a default. I also was told once that running FDS with
single precision might actually cost more in CPU time because of the
need to truncate native 8 byte reals to 4 byte. I suspect that the
GPUs are designed for 4 byte reals.

You have probably figured out that you can reduce the real precision
by changing the definition of "EB" in prec.f90.

> > > Hendrik- Hide quoted text -
>
> - Show quoted text -

[D Perron]

This is very interesting. Is there any way to get a CUDA compatible
version of FDS, or, at least a detailed set of instructions as to the
changes needed in the FDS code?

I have been looking at using the CUDA based systems for other CFD
programs and it looks to be very promising, namely due to the
relatively small maximum FLOP limit of desktop level machines with the
current x32 or x64 structure.

Cheers,
David

[Hendrik]

You're right. The peek performance of the GTX 280 I am using is 933
GFlops in single presicion mode and only 117 GFlops using double
prescision. As you may noticed, the performance of the single
presicion version in my sample files is only twice as high as the
double presicion version (and not 8 times). Why is that?

I do belief the bottelneck just now is the memory access pattern.
Since the GPU needs to access the memory in blocks an uncoordinated
reading or writing from or to the memory (as done in many loops in FDS
at present) forces many processors to switch to an idle mode after
finishing one computation. I mean to say: in my currently version not
all of the 240 processors are in use at any time which reduces the
maximum performance significantly.