GPU Geek's blog

CUDA DSMC paper has been published

2012-01-24T23:09:00.000-06:00

My CUDA DSMC paper has finally been published in SIMULATION journal. You can find it here.

Thanks!

Dropped out

2011-11-09T14:00:00.000-06:00

I dropped out from the school and working at a start up company in San Francisco now. I'll continue blogging soon...

Call me from the blog!

2011-05-11T15:55:00.000-05:00

Thanks to my friends from Zingaya, Ltd you now can call me from this blog by clicking the button on the right side (Call me!). To do so, you just need to have a Flash player and microphone installed. You can call from anywhere and the call will be redirected to my cell phone.

Amazing technology indeed!

SpringSim 2011

2011-04-08T13:10:00.000-05:00

I went to Boston for SpringSim 2011 this week for presenting GPU Smoldyn. It was fun. I liked the city a lot. I also met lots of interesting people there. You can check out the Smoldyn paper I was presenting and the presentation itself at the following links:

Paper
Slides

Random number generators on GPU

2011-03-03T17:04:00.000-06:00

Almost every simulation requires random numbers generation (RNG) for achieving accurate results. In my simulations I used 2 RNGs. MersenneTwister - has a very long period and can be used in simulations that are supposed to be running for a long time, and XOR128 - extremely fast RNG, however its period is not as large as MerseneTwister's one. I created a small library that contains 3 RNGs: MersenneTwister, XOR128 and Tausworthe generator.

For implementing MersenneTwister I deeply modified NVIDIA's sample. Each thread in a grid runs its own RNG. States are kept in a global memory. Each thread is supposed to read a state in the beginning of execution and save it at the end. I have a helper classes, which do it.
The library is available at link. And it contains an example file, however I'm going to put some small tutorial here.

The library consists of 2 parts - device and host.

Device code for RNGs is defined in files:

mtwister.cuh - for MersenneTwister
rndfast.cuh - for XOR128 and Tausworthe

So if you want to use these generators in your projects you just need to include one of these files in your CU file. Then in kernels you can use generators like this (MersenneTwister example):

const int maxValue = 1024;
const int tid = blockDim.x*blockIdx.x + 
                 threadIdx.x;

MTGenerator gen(tid);

int   intNum   = gen.GetInt(maxValue);
flaot floatNum = gen.GetFloat();

If you want to use other generators just use classes TausGenerator for Tausworthe and Xor128Generator for XOR128.

You will also need to initialize seeds. These functions are defined in files:

mtwister.h - for MersennwTwister
seedgen.h - for XOR128 and Tausworthe

Simply call one of these functions for initialization:

void InitGPUTwisters(const char* fname, 
                     int numGen, int seed);
void InitFastRngs(int numRngs);

And one of these functions for deinitialization:

void DeinitGPUTwisters();
void DeinitFastRngs();

Again, there is a good example in the package, so feel free to download and use it. If you have any questions, please send me an email or comment.

Web access to GPU Smoldyn

2011-03-01T23:42:00.000-06:00

In my previous post I described a GPU implementation of Smoldyn. In this post I'm going to describe a Web-Service that provides an internet access to it. Using this service a user can start a simulation just uploading a configuration file describing a model. The service executes this file and uploads screenshots and other output files to user's profile, so a user can access them. This allows users, who doesn't have access to high performance GPUs enjoy all benefits of GPU implementation (200x speedup for instance).

The service consists of 2 parts. The first part is a web-interface, that manages user's profile and provides access to information is implemented using django framework. The database it uses is a MySql database. The second part is an application that executes simulations themselves. Currently it is a python script that fetches data from the database, executes simulations using this data (the script executes an original GPU Smoldyn application written in CUDA C++ using OS services). Then the script uploads the output information to the database. Current architecture is presented on the image below.

In the nearest future I'm going to improve this service by adding Amazon Simple Storage (SSS) support for storage user files like output files and screenshots. Also I'm going to add a new level - dispatcher level that would be providing jobs for execution nodes, achieving better scalability. The future architecture is presented on the picture below.

You can try the service by clicking the link. You would need to register, but it is free and there is no any obligations.

Biochemical simulations on GPU

2011-02-23T18:19:00.006-06:00

My second project is a CUDA implementation of Smoldyn simulator. Smoldyn is cell-scale biochemical simulator which simulates each molecule of interest individually to capture natural stochasticity and for nanometer-scale spatial resolution. This is a particle-based method that allows simulation of chemical molecules with Brownian motion and different reaction types.

The current version of GPU Smoldyn supports the following features:

Isotropic diffusion;
Zero/First/Second order reactions;
User commands;
Surfaces interactions;

So, using this functionality different molecule to surface interactions can be simulated. Reaction-diffusion systems, like a molecular mutation (A->B), synthesis of molecules, or even bi-molecular reactions (A+B->C) can be simulated as well.

Implementation details

This simulation is implemented for GPU using CUDA. It is implemented for Linux, so OpenGL was used for visualization.

I widely used Thrust in this simulation, especially for operations like stream compaction, reduction and scan. For second order reaction processing I needed sorting for spatial hashing. Firstly I used thrust::sort, however, it turned out, that NVIDIA's radix sort kernel is much faster than thrust::sort, because in NVIDIA's implementation you can specify maximum number of sorting bits, which is very convenient for cases, when the maximum value for a sorting key is known. In my case, it is a number of cells-1.

The tricky aspect of this simulation is a dynamic memory allocation, since molecules can die and re spawn during the simulation. Every molecule has a type field. In fact, a molecule is described as:

struct molecule_t
{
float3 pos;
int type;
};

Type equal to -1 represents dead molecule. I allocate enough memory to handle new particles creation during the simulation. At the initialization stage the molecule array is filled with a desired number of molecules, and the rest of it is initialized with dead molecules (with type == -1). A pointer to the last alive molecule is maintained in a global memory. If during the simulation I need to allocate a molecule, the pointer is advanced using atomicAdd and new molecule is generated (new molecules are spawned at the end of the array, maintaining proper order of alive molecules). If I need to deallocate molecule, I just set its type to -1 and then do compaction after the kernel finishes. I use thrust::remove_if operation to move dead molecules to the end of the array.

I achieved average speedup of 200x for this implementation against original Smoldyn and this implementation is statistically correct. I wrote a paper about this implementation and I'll go to SpringSim 2011 to present it, so if you're interested, you can catch me there (or send me an email). I'll post the paper here after it is published in the proceedings.

Some screenshots

GPU-based Direct Simulation Monte Carlo

2011-02-21T21:56:00.002-06:00

It's been a while since my last post. Lots of things have changed. Now I'm a PhD student at UW-Milwaukee and a research assistant at the Complex Systems Simulation Lab. In this post I will describe my first project which is a GPU (CUDA) implementation of DSMC method for particle simulations.

DSMC is a computational method for fluid mechanics simulation. Simply put, it can be used for simulation of interaction between gas molecules and a solid body. For example, the case, when a space ship enters the atmosphere can be simulated. The method is relatively simple, but it requires simulation of huge number of molecules for an accurate result, so it is a good candidate for parallelization.

Implementation details

Since it is a Monte Carlo-based method, it requires a random number generator. The Mersenne Twister is one of the best RNGs at the moment, however it has a large state which doesn't make it very suitable for GPU, however this RNG was used in the simulation. I modified the NVIDIA's sample, so each thread in the simulation can run its own twister. Also XOR128 was implemented as well and is used for very large simulations.

For inter-particle collisions I used Bird's collision method, which is currently the most accurate one. Since it is a particle-based simulation, I adopted particles' hashing method from NVIDIA's particles sample (using radixsort).

During the simulation particles can collide with an obstacle (polygonal object). For more efficient processing this object is subdivided using a regular grid. For AABB to triangle testing I used very fast and accurate method by Moller. There is also very nice ray to triangle testing method by the same author, I used this method as well. The simulation can easily run in a real time with systems up to 16 mln particles and with a complex polygonal object in it.

As a result, I achieved a speedup of 65x over a CPU implementation and got some nice screenshots:

In these screenshots a small window in a top right corner shows per-cell concentration. Red is high, blue is low, green is in between. The blue cells in the simulation show the current grid slice that is presented in the concentration window.

I also submitted a paper to SIMULATION journal about this method. It got accepted, so I'll post it here when the journal is released. I will also describe particular aspects of this implementation in more details in my further posts, since it will take lots of space.

This was my first project, but I have some more to describe, so stay tuned :)

VNC for Mac

2009-05-25T07:04:00.000-05:00

Set up remote access to my Mac. Server is VineVNC for Mac OS X. Client is TightVNC for WinXP. Works well. Now I can develop iPhone applications from around the world :) (even from iPhone with Mocha VNC).

Java bytecode cross-compilation for iPhone

2009-05-24T06:01:00.000-05:00

Just have seen a pretty interesting lecture about Java bytecode cross compilation for iPhone on Google Tech Talks. I assume this may let iPhone developers not to study ObjC (which pretty nice though). I guess it would be interesting to try load such cross-compiled application to real device or even to AppStore. Maybe I will try to do that.

Lecture

Slides

Check it out later :)

Hello World!

2008-10-08T05:15:00.000-05:00

Hi, here is a first message in my shiny new blog :)

Check it out later!