So I've got a decent head for what problems I work with are best one in serial, and which can be managed in parallel. But right now, I don't have much of an idea of what's best handled by CPU-based computation, and what should be offloaded to a GPU.
I know its a basic question, but much of my searching gets caught in people clearly advocating for one or the other without really justifying why, or somewhat vague rules of thumb. Looking for a more useful response here.
GPU hardware has two particular strengths: raw compute (FLOPs) and memory bandwidth. Most difficult computational problems fall into one of these two categories. For example, dense linear algebra (A * B = C or Solve[Ax = y] or Diagonalize[A], etc) falls somewhere on the compute/memory bandwidth spectrum depending on system size. Fast Fourier transforms (FFT) also fit this mold with high aggregate bandwidth needs. As do other transformations, grid/mesh-based algorithms, Monte Carlo, etc. If you look at the NVIDIA SDK code examples, you can get a feel for the sorts of problems that are most commonly addressed.
I think the more instructive answer is to the question `What kinds of problems are GPUs really bad at?' Most problems that don't fall into this category can be made to run on the GPU, though some take more effort than others.
Problems that don't map well are generally too small or too unpredictable. Very small problems lack the parallelism needed to use all the threads on the GPU and/or could fit into a low-level cache on the CPU, substantially boosting CPU performance. Unpredictable problems have too many meaningful branches, which can prevent data from efficiently streaming from GPU memory to the cores or reduce parallelism by breaking the SIMD paradigm (see 'divergent warps'). Examples of these kinds of problems include:
__synchtreads()).
$\endgroup$
Problems which have a high arithmetic intensity and regular memory access patterns are typically easy(ier) to implement on GPUs, and perform well on them.
The basic difficulty in having high performance GPU code is that you have a ton of cores, and you want them to all be utilized to their full potency as much as possible. Problems which have irregular memory access patterns or do not have high arithmetic intensity make this difficult: either you spend a long time communicating results or you spend a long time fetching stuff from memory (which is slow!), and not enough time crunching numbers. Of course the potential for concurrency in your code is critical to its ability to be implemented well on GPU as well.
This is not intended as an answer on its own but rather an addition to the other answers by maxhutch and Reid.Atcheson.
To get the best out of GPUs your problem does not only need to be highly (or massively) parallel, but also the core algorithm that will be executed on the GPU, should be as small as possible. In OpenCL terms this is mostly referred as the kernel.
To be more precise, the kernel should fit into the register of each multiprocessing unit (or compute unit) of the GPU. The exact size of the register is dependent on the GPU.
Given the kernel is small enough, the raw data of the problem needs to fit into the GPU's local memory (read: local memory (OpenCL) or shared memory (CUDA) of a compute unit). Otherwise even the high memory bandwidth of the GPU is not fast enough to keep the processing elements busy all the time.
Usually this memory is about 16 to 32 KiByte big.
Probably a more technical addition to the previous replies: CUDA (i.e. Nvidia) GPUs can be described as a set of processors that work autonomously on 32 threads each. The threads in each processor work in lock-step (think SIMD with vectors of length 32).
Although the most tempting way to work with GPUs is to pretend that absolutely everything runs in lock-step, this is not always the most efficient way of doing things.
If your code does not parallelize nicely/automatically to hundreds/thousands of threads, you may be able to break it down into individual asynchronous tasks that do parallelize well, and execute those with only 32 threads running in lock-step. CUDA provides a set of atomic instructions which make it possible to implement mutexes which in turn allows the processors to synchronize among themselves and process a list of tasks in a thread pool paradigm. Your code would then work much in the same way as it does on a multi-core system, just keep in mind that each core then has 32 threads of its own.
Here's a small example, using CUDA, of how this works
/* Global index of the next available task, assume this has been set to
zero before spawning the kernel. */
__device__ int next_task;
/* We will use this value as our mutex variable. Assume it has been set to
zero before spawning the kernel. */
__device__ int tasks_mutex;
/* Mutex routines using atomic compare-and-set. */
__device__ inline void cuda_mutex_lock ( int *m ) {
while ( atomicCAS( m , 0 , 1 ) != 0 );
}
__device__ inline void cuda_mutex_unlock ( int *m ) {
atomicExch( m , 0 );
}
__device__ void task_do ( struct task *t ) {
/* Do whatever needs to be done for the task t using the 32 threads of
a single warp. */
}
__global__ void main ( struct task *tasks , int nr_tasks ) {
__shared__ task_id;
/* Main task loop... */
while ( next_task < nr_tasks ) {
/* The first thread in this block is responsible for picking-up a task. */
if ( threadIdx.x == 0 ) {
/* Get a hold of the task mutex. */
cuda_mutex_lock( &tasks_mutex );
/* Store the next task in the shared task_id variable so that all
threads in this warp can see it. */
task_id = next_task;
/* Increase the task counter. */
next_tast += 1;
/* Make sure those last two writes to local and global memory can
be seen by everybody. */
__threadfence();
/* Unlock the task mutex. */
cuda_mutex_unlock( &tasks_mutex );
}
/* As of here, all threads in this warp are back in sync, so if we
got a valid task, perform it. */
if ( task_id < nr_tasks )
task_do( &tasks[ task_id ] );
} /* main loop. */
}
You then have to call the kernel with main<<<N,32>>>(tasks,nr_tasks) to make sure that each block contains only 32 threads and thus fits in a single warp. In this example I also assumed, for simplicity, that the tasks do not have any dependencies (e.g. one task depends on the results of another) or conflicts (e.g. work on the same global memory). If this is the case, then the task selection becomes a bit more complicated, but the structure is essentially the same.
This is, of course, more complicated than just doing everything on one large batch of cells, but significantly broadens the type of problems for which GPUs can be used.
One point not made so far is that the current generation of GPU's don't do as well at double precision floating point computations as with single precision computations. If your computations have to be done in double precision, then you can expect the run time to increase by a factor of 10 or so over single precision.
From a metaphorical point of view, the gpu can be seen as a person lying on a bed of nails. The person lying on top is the data and in the base of each nail there is a processor, so the nail is actually an arrow pointing from processor to memory. All nails are in a regular pattern, like a grid. If the body is well spread, it feels good (performance is good), if the body only touches some spots of the nail bed, then the pain is bad (bad performance).
This can be taken as a complementary answer to the excellent answers above.
Old question, but I think that this answer from 2014 - related to statistical methods, but generalisable for anyone who knows what a loop is - is particularly illustrative and informative.
GPUs have long latency I/O, so lots of threads need to be used to saturate the memory. To keep a warp busy requires lots of threads. If the code path is 10 clocks and I/O latency 320 clocks, 32 threads should come close to saturating the warp. If the code path is 5 clocks, then double the threads.
With a thousand cores, look for thousands of threads to fully utilize the GPU.
Memory access is by cache line, usually 32 bytes. Loading one byte has comparable cost to 32 bytes. So, coalesce the storage to increase locality of usage.
There are lots of registers and local RAM to each warp, allowing for neighbor sharing.
Proximity simulations of large sets should optimize well.
Random I/O and single threading is a kill joy...
Imagine a problem that can be solved by lots of brute force, like Travelling Salesman. Then imagine you've got racks of servers with 8 spanky video cards each, and each card has 3000 CUDA cores.
Simply solve ALL the possible salesman's routes and then sort for time/distance/some metric. Sure you're throwing away almost 100% of your work, but brute force is a viable solution sometimes.
From studying many Engineering ideas, I woulds say a gpu is a form of focusing of tasks,of memory-management,of repeatable-calculation.
Many formulas might be simple to write but painful to calculate such as in matrix mathematics you don't get a single answer but many values.
This is important in computing as how fast a computer is calculating values and running formulas as some formulas can't run without all calculated values (hence slow down). A computer doesn't know very well what order to run formulas or calculate values to use in these programs. It mainly brute forces through at fast speeds and breaks formulas into chucks to calculate but many programs these days require these calculated chucks right now and wait in ques (and ques of ques and more ques of ques).
For example in a simulation game which should be calculated first in collisions the damage of the collision, the position of the objects, the new velocity? How much time should this take? How can any cpu handle this load? Also, most programs are very abstract requiring more time to handle data and aren't always designed for multi-threading or aren't any good ways in abstract programs to do this effectively.
As cpu became better and better people became sloppy in programming and we must program for many different types of computers as well. A gpu is designed to brute force through many simple calculations at the same time (not mention memory (secondary/ram) and heating cooling are the main bottle necks in computing). A cpu is managing many many ques at the same time or being pulled in to many directions it is figuring out what to do not being able to do it. (hey it's almost human)
A gpu is grunt worker the tedious work. A cpu is managing complete chaos and can't handle every detail.
So what do we learn? A gpu does detail tedious work all at once and a cpu is a multi-task machine that can't focus very well with too many tasks to do. (It's like it has attention disorder and autism at the same time).
Engineering there is the ideas, design, reality, and a lot of grunt work.
As I leave remember to start simple, start quickly, fail-quickly, fail-fast, and never stop trying.
