# What kinds of problems lend themselves well to GPU computing?

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:

• Most graph algorithms (too unpredictable, especially in memory-space)
• Sparse linear algebra (but this is bad on the CPU too)
• Small signal processing problems (FFTs smaller than 1000 points, for example)
• Search
• Sort
• Still, GPU solutions for those "unpredictable" problems are possible and, while nowadays not typically feasible, may gain significance in the future. – leftaroundabout Jan 23 '12 at 11:18
• I'd like to specifically add branches to the list of GPU performance breakers. You want all your (hundreds) execute the same instruction (as in SIMD) to perform truly parallel computation. For example, on AMD cards if any of the instruction flows encounters a branch and must diverge - all the wavefront (parallel group) diverges. If another units from the wavefront must not diverge - they must perform a second pass. That's what maxhutch means by predictability, I guess. – Violet Giraffe Jan 23 '12 at 14:13
• @VioletGiraffe, that's not necessarily true. In CUDA (i.e. on Nvidia GPUs), branch divergence only affects the current warp, which is at most 32 threads. Different warps, although executing the same code, are not synchronous unless explicitly synchronized (e.g. with __synchtreads()). – Pedro Jan 23 '12 at 23:33
• @Pedro: True, but branching in general does hurt performance. For high-performance codes (what GPU code isn't?), it is almost essential to take that into account. – jvriesem Aug 11 '15 at 1:20

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.

• Can you specify what you mean by regular memory access patterns? – Fomite Jan 23 '12 at 6:05
• maxhutch's answer is better than mine. What I mean by regular access pattern is that memory is accessed in a temporal and spatially locally way. That is: you do not make huge jumps around memory repeatedly. It's also something of a package deal I've noticed. It also is taken to mean that your data access patterns can be predetermined either by the compiler somehow or by you the programmer so that branching (conditional statements in code) is minimized. – Reid.Atcheson Jan 23 '12 at 6:14

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.

• Isn't the local/shared memory of each processing unit shared among all of the dozens(?) of threads running within a single cluster of cores? In this case, don't you actually need to keep your working set of data significantly smaller in order to get full performance out of the GPU? – Dan Neely Jan 23 '12 at 13:39
• The local/shared memory of a processing unit is only accessible by the compute unit itself and thus only shared by the processing elements of this compute unit. The global memory of the graphics card (usually 1GB) is accessible by all processing units. The bandwidth between the processing elements and the local/shared memory is very fast (>1TB/s) but the bandwidth to the global memory a lot slower (~100GB/s) and needs to be shared among all compute units. – Torbjörn Jan 23 '12 at 13:43
• I wasn't asking about the main GPU memory. I thought the on die memory was only allocated at the cluster of core level not per individual core. ex for an nVidia GF100/110 gpu; for each of the 16 SM clusters not the 512 cuda cores. With each SM designed to run up to 32 threads in parallel maximizing GPU performance would require keeping the working set in the 1kb/thread range. – Dan Neely Jan 23 '12 at 14:02
• @Torbjoern What you want is to keep all the GPU execution pipelines busy, GPUs achieve this two ways: (1) the most common way is to increase occupancy, or said differently, by increasing the number of concurrent threads (small kernels use less of the shared resources so you can have more active threads); maybe better, is to (2) increase the instruction level parallelism within your kernel, so you can have larger kernel with relatively low occupancy (small number of active threads). See bit.ly/Q3KdI0 – fcruz Aug 6 '12 at 14:21

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. */

/* We will use this value as our mutex variable. Assume it has been set to
zero before spawning the kernel. */

/* 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 );
}

/* Do whatever needs to be done for the task t using the 32 threads of
a single warp. */
}

/* The first thread in this block is responsible for picking-up a task. */
if ( threadIdx.x == 0 ) {

/* Get a hold of the task mutex. */

/* Store the next task in the shared task_id variable so that all
threads in this warp can see it. */

/* Increase the task counter. */
next_tast += 1;

/* Make sure those last two writes to local and global memory can
be seen by everybody. */

/* Unlock the task mutex. */

}

/* As of here, all threads in this warp are back in sync, so if we
got a valid task, perform it. */

} /* 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.

• This is technically true, but high parallelism is needed to get high memory bandwidth and there is a limit to the number of asynchronous kernel calls (currently 16). Thee is also tons of undocumented behavior related to scheduling in the current release. I'd advise against relying on asynchronous kernels to imporove performance for the time being... – Max Hutchinson Jan 23 '12 at 17:16
• What I'm describing can be done all in one single kernel call. You can make N blocks of 32 threads each, such that each block fits into a single warp. Each block then acquires a task from a global task list (access controlled using atomics/mutexes) and computes it using 32 lock-stepped threads. All this happens in a single kernel call. If you'd like a code example, let me know and I'll post one. – Pedro Jan 23 '12 at 22:51

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.

• I want to disagree. Most (or all) newer GPU's have native double precision support. Almost every such GPU reports double precision computations running at approximately half the speed of single precision, likely due to the simple doubling of required memory accesses/bandwidth. – Godric Seer Jun 7 '13 at 16:15
• While it's true that the latest and greatest Nvidia Tesla cards do offer peak double precision performance that's half of the peak single precision performance, the ratio is 8 to 1 for more common Fermi architecture consumer grade cards. – Brian Borchers Jun 8 '13 at 1:17
• @GodricSeer The 2:1 ratio of SP and DP floating-point has very little to do with bandwidth and almost everything to do with how many hardware units exist to execute these operations. It is common to reuse the register file for SP and DP, hence the floating-point unit can execute 2x the SP ops as DP ops. There are numerous exceptions to this design, e.g. IBM Blue Gene/Q (does not have SP logic and thus SP runs at ~1.05x DP). Some GPUs have ratios other than 2, e.g. 3 and 5. – Jeff Jan 3 '18 at 20:57
• It's four years since I wrote this answer, and the current situation with NVIDIA GPU's is that for the GeForce and Quadro lines, the DP/SP ratio is now 1/32. NVIDIA's Tesla GPUs have much stronger double precision performance but also cost a lot more. On the other hand, AMD hasn't crippled double precision performance on its Radeon GPU's in the same way. – Brian Borchers Jan 3 '18 at 23:14

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...

• This is a genuinely fascinating question; I am arguing with myself as to whether it's possible (or worth the effort) to 'parallel-ise' a reasonably straightforward task (edge detection in aerial images) when each task takes ~0.06sec but there are ~1.8 million tasks to perform (per year, for 6 years' worth of data: the tasks are definitely separable)... thus ~7.5 days' worth of compute time on one core. If each calc was faster on a GPU, and the job could be parallelised 1-per-nGPUcores [n small], is it actually likely that the job time could drop to ~1 hour? Seems unlikely. – GT. May 16 '15 at 7:13

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.

• I had access to a small farm of 4 such servers for a week, and in five days I did more distributed.net blocks than in the previous 10 years. – Criggie Oct 31 '16 at 8:14

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.