CUDA: How many concurrent threads in total?

by

The GTX 580 can have 16 * 48 concurrent warps (32 threads each) running at a time. That is 16 multiprocessors (SMs) * 48 resident warps per SM * 32 threads per warp = 24,576 threads.

Don’t confuse concurrency and throughput. The number above is the maximum number of threads whose resources can be stored on-chip simultaneously — the number that can be resident. In CUDA terms we also call this maximum occupancy. The hardware switches between warps constantly to help cover or “hide” the (large) latency of memory accesses as well as the (small) latency of arithmetic pipelines.

While each SM can have 48 resident warps, it can only issue instructions from a small number (on average between 1 and 2 for GTX 580, but it depends on the program instruction mix) of warps at each clock cycle.

So you are probably better off comparing throughput, which is determined by the available execution units and how the hardware is capable of performing multi-issue. On GTX580, there are 512 FMA execution units, but also integer units, special function units, memory instruction units, etc, which can be dual-issued to (i.e. issue independent instructions from 2 warps simultaneously) in various combinations.

Taking into account all of the above is too difficult, though, so most people compare on two metrics:

  1. Peak GFLOP/s (which for GTX 580 is 512 FMA units * 2 flops per FMA * 1544e6 cycles/second = 1581.1 GFLOP/s (single precision))
  2. Measured throughput on the application you are interested in.

The most important comparison is always measured wall-clock time on a real application.

More Related Contents:

  1. Modifying registry to increase GPU timeout, windows 7
  2. How to measure the inner kernel time in NVIDIA CUDA?
  3. Passing Host Function as a function pointer in __global__ OR __device__ function in CUDA
  4. How do CUDA blocks/warps/threads map onto CUDA cores?
  5. How are 2D / 3D CUDA blocks divided into warps?
  6. nvidia-smi Volatile GPU-Utilization explanation?
  7. How do I use Nvidia Multi-process Service (MPS) to run multiple non-MPI CUDA applications?
  8. CUDA apps time out & fail after several seconds – how to work around this?
  9. How does CUDA assign device IDs to GPUs?
  10. Unspecified launch failure on Memcpy
  11. Different CUDA versions shown by nvcc and NVIDIA-smi
  12. Thrust inside user written kernels
  13. What is the purpose of using multiple “arch” flags in Nvidia’s NVCC compiler?
  14. What can I do against ‘CUDA driver version is insufficient for CUDA runtime version’?
  15. What is a bank conflict? (Doing Cuda/OpenCL programming)
  16. Can anyone provide sample code demonstrating the use of 16 bit floating point in cuda?
  17. Any particular function to initialize GPU other than the first cudaMalloc call?
  18. CUDA and nvcc: using the preprocessor to choose between float or double
  19. Polymorphism and derived classes in CUDA / CUDA Thrust
  20. Why is NVIDIA Pascal GPUs slow on running CUDA Kernels when using cudaMallocManaged
  21. How do I select which GPU to run a job on?
  22. multi-GPU basic usage
  23. How can I make tensorflow run on a GPU with capability 2.x?
  24. How to implement device side CUDA virtual functions?
  25. How to create a CUDA context?
  26. Why does my CUDA kernel crash (unspecified launch failure) with a different dataset size?
  27. CUDA determining threads per block, blocks per grid
  28. How to disable a specific nvcc compiler warnings
  29. CUDA function pointers
  30. Using an array of device function pointers

Leave a Comment