GPU Programming

Learn the principles of GPU programming with multiple-choice questions: the grid/block/warp hierarchy, shared memory, coalesced memory access, warp divergence and occupancy.

—Your accuracy
0 / 50Questions answered
Practice this topicAll questions in this topic, in random order Weak questionsRetry only the questions you last got wrong

GPU Programming: key points

A GPU is a processor that excels at running thousands to tens of thousands of threads simultaneously, applying the same operation to large amounts of data. Programs express parallelism in a hierarchy of grids, blocks and threads, and the hardware executes groups of about 32 threads, called warps, in lockstep.

Getting performance requires coalesced memory access, where neighboring threads read neighboring addresses; on-chip shared memory used within a block; avoiding divergence within a warp; and enough occupancy to hide memory latency. This topic focuses on principles that do not depend on product generations.

All GPU Programming questions (50)

  1. How do the design philosophies of CPUs and GPUs differ?
  2. Which characteristic best describes work that GPUs process quickly?
  3. Which processing tends to run inefficiently on a GPU?
  4. In GPU programming such as CUDA, what is a "kernel"?
  5. Which correctly orders the CUDA thread hierarchy from largest to smallest?
  6. A CUDA kernel computes each thread's index i with the expression above. What is i for the thread with blockIdx.x = 2, blockDim.x = 256 and threadIdx.x = 5?
  7. To process a 1000-element array with one element per thread and 256 threads per block, how many blocks should be launched?
  8. Why does the CUDA kernel above include "if (i < n)"?
  9. On NVIDIA GPUs, what is the unit of 32 threads bundled together to execute the same instruction in lockstep called?
  10. How does the GPU SIMT (single instruction, multiple threads) model differ from CPU SIMD instructions?
  11. Within one warp, a conditional branch sends half the threads to the if path and half to the else path (warp divergence). How is this generally executed?
  12. Both (1) and (2) above make half the threads run A and the other half run B. Which is correct about warp efficiency (32 threads per warp)?
  13. Which describes memory coalescing on a GPU?
  14. With i as each thread's index, reading the array every 32 elements as above is slower than reading a[i]. What is the main reason?
  15. When processing particle data with positions x, y and z on a GPU, which layout is better for memory coalescing?
  16. What is a characteristic of CUDA shared memory?
  17. What does CUDA's __syncthreads() do?
  18. In CUDA, __syncthreads() is placed inside "if (threadIdx.x < 16)" so that only some threads reach it. What can happen?
  19. Which describes a bank conflict in CUDA shared memory?
  20. A 32 × 32 float tile in shared memory causes bank conflicts when a warp's 32 threads read down a column. What is the common fix?
  21. When multiplying large matrices on a GPU, why are the matrices split into small tiles loaded into shared memory?
  22. What can happen if a GPU kernel uses too many registers per thread?
  23. Which describes occupancy on a GPU?
  24. What is the main way a GPU hides long memory latency (hundreds of cycles)?
  25. Which statement about GPU occupancy is correct?
  26. In CUDA, what is the relationship between blocks and SMs (streaming multiprocessors)?
  27. In CUDA, you want to wait mid-kernel for every other block to finish before moving to the next phase. What is the usual approach?
  28. Building a histogram of pixel values on a GPU can be slow if every thread does atomicAdd directly on 256 bins in global memory. What is a good improvement?
  29. When summing many values on a GPU, why use a tree-shaped reduction (repeatedly adding neighbors to halve the count)?
  30. How many stages does a tree-shaped reduction need to sum 1024 values?
  31. In a program using a discrete GPU, what is generally true about CPU-side and GPU-side memory?
  32. Sending data to a GPU for light work such as adding array elements once can be slower than computing on the CPU. What is the main reason?
  33. Why does pinning (page-locking) CPU memory before transferring it to the GPU speed up the transfer?
  34. What is the main purpose of using multiple CUDA streams?
  35. A program that launches a very small kernel thousands of times in a row is slow. What is the most suitable improvement?
  36. Which describes CUDA unified memory?
  37. When a GPU performs element-wise addition of two large arrays (c[i] = a[i] + b[i]), what mainly determines performance?
  38. Why is double-precision (FP64) performance on many consumer gaming GPUs much lower than single-precision (FP32), for example one thirty-second?
  39. What is a characteristic of the tensor cores (matrix units) in recent GPUs?
  40. In GPU mixed-precision computation, why are multiplications often done in a low precision such as FP16 while accumulation is done in FP32?
  41. Summing the same floating-point data on a GPU gives results whose last digits differ slightly from run to run. What is the most likely cause?
  42. What is a common guideline for choosing the number of threads per block in CUDA?
  43. What is the advantage of writing a kernel with a loop like the one above (a grid-stride loop)?
  44. For what kind of access does CUDA constant memory perform especially well?
  45. In general, how do GPU caches compare with CPU caches?
  46. Which is an example of a parallel-programming framework for GPUs that also works on GPUs other than NVIDIA's?
  47. On a GPU you want to compact only the elements meeting a condition into another array (each thread outputs 0 or 1 items). Which parallel computation is commonly used to determine each thread's write position?
  48. Which assumption must be respected when writing GPU kernels?
  49. When transposing a large matrix on a GPU (b[j][i] = a[i][j]), what is the main purpose of going through shared memory?
  50. In CUDA, what generally happens right after the CPU-side program launches a kernel?

Other topics in Computer systems