How do I choose grid and block dimensions for CUDA kernels?

by

There are two parts to that answer (I wrote it). One part is easy to quantify, the other is more empirical.

Hardware Constraints:

This is the easy to quantify part. Appendix F of the current CUDA programming guide lists a number of hard limits which limit how many threads per block a kernel launch can have. If you exceed any of these, your kernel will never run. They can be roughly summarized as:

  1. Each block cannot have more than 512/1024 threads in total (Compute Capability 1.x or 2.x and later respectively)
  2. The maximum dimensions of each block are limited to
    [512,512,64]/[1024,1024,64] (Compute 1.x/2.x or later)
  3. Each block cannot consume more than 8k/16k/32k/64k/32k/64k/32k/64k/32k/64k registers total
    (Compute 1.0,1.1/1.2,1.3/2.x-/3.0/3.2/3.5-5.2/5.3/6-6.1/6.2/7.0)
  4. Each block cannot consume more than 16kb/48kb/96kb of shared memory (Compute
    1.x/2.x-6.2/7.0)

If you stay within those limits, any kernel you can successfully compile will launch without error.

Performance Tuning:

This is the empirical part. The number of threads per block you choose within the hardware constraints outlined above can and does effect the performance of code running on the hardware. How each code behaves will be different and the only real way to quantify it is by careful benchmarking and profiling. But again, very roughly summarized:

  1. The number of threads per block should be a round multiple of the warp size, which is 32 on all current hardware.
  2. Each streaming multiprocessor unit on the GPU must have enough active warps to sufficiently hide all of the different memory and instruction pipeline latency of the architecture and achieve maximum throughput. The orthodox approach here is to try achieving optimal hardware occupancy (what Roger Dahl’s answer is referring to).

The second point is a huge topic which I doubt anyone is going to try and cover it in a single StackOverflow answer. There are people writing PhD theses around the quantitative analysis of aspects of the problem (see this presentation by Vasily Volkov from UC Berkley and this paper by Henry Wong from the University of Toronto for examples of how complex the question really is).

At the entry level, you should mostly be aware that the block size you choose (within the range of legal block sizes defined by the constraints above) can and does have a impact on how fast your code will run, but it depends on the hardware you have and the code you are running. By benchmarking, you will probably find that most non-trivial code has a “sweet spot” in the 128-512 threads per block range, but it will require some analysis on your part to find where that is. The good news is that because you are working in multiples of the warp size, the search space is very finite and the best configuration for a given piece of code relatively easy to find.

More Related Contents:

  1. Performance optimization strategies of last resort [closed]
  2. What is the best way to set a register to zero in x86 assembly: xor, mov or and?
  3. Why are loops always compiled into “do…while” style (tail jump)?
  4. One could use a profiler, but why not just halt the program? [closed]
  5. Google app script timeout ~ 5 minutes?
  6. Google app script timeout ~ 5 minutes?
  7. How are x86 uops scheduled, exactly?
  8. What methods can be used to efficiently extend instruction length on modern x86?
  9. What setup does REP do?
  10. When, if ever, is loop unrolling still useful?
  11. Is ADD 1 really faster than INC ? x86 [duplicate]
  12. How is CUDA memory managed?
  13. Why is a conditional move not vulnerable to Branch Prediction Failure?
  14. How to measure the inner kernel time in NVIDIA CUDA?
  15. Can modern x86 implementations store-forward from more than one prior store?
  16. What is the most ridiculous pessimization you’ve seen? [closed]
  17. What are the major performance hitters in AS3 aside from rendering vectors?
  18. nvidia-smi Volatile GPU-Utilization explanation?
  19. How do I use Nvidia Multi-process Service (MPS) to run multiple non-MPI CUDA applications?
  20. Unexpectedly poor and weirdly bimodal performance for store loop on Intel Skylake
  21. Why do these goroutines not scale their performance from more concurrent executions?
  22. Best Practices for Multiple OnEdit Functions
  23. Memory Allocation/Deallocation Bottleneck?
  24. why do we need cudaDeviceSynchronize(); in kernels with device-printf?
  25. Relative performance of x86 inc vs. add instruction
  26. How does CUDA assign device IDs to GPUs?
  27. How can the rep stosb instruction execute faster than the equivalent loop?
  28. JMeter : How to record HTTPS traffic?
  29. Fastest way to strip all non-printable characters from a Java String
  30. Return address prediction stack buffer vs stack-stored return address?

Leave a Comment