GPU Programming

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?

  1. ACall __syncthreads()
  2. BSplit the kernel at that point and launch the second kernel after the first finishes
  3. CSpin in a loop waiting for flags written by other blocks
Show answer and explanation

Answer: B. Split the kernel at that point and launch the second kernel after the first finishes

Blocks are independent and may run in any order, so not all of them are necessarily resident on the GPU at once. If the block being waited for has not started, a flag-waiting loop never ends. The usual approach relies on the fact that kernels launched in the same stream start only after the previous one finishes, splitting the kernel to obtain a global synchronization.

  • A. __syncthreads() synchronizes only within a block.
  • B. Correct. The kernel boundary acts as a global synchronization.
  • C. There is no guarantee all blocks are resident, so it can hang.
Practice this topic Back to topic

More GPU Programming questions

  1. 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?
  2. When summing many values on a GPU, why use a tree-shaped reduction (repeatedly adding neighbors to halve the count)?
  3. How many stages does a tree-shaped reduction need to sum 1024 values?
  4. In a program using a discrete GPU, what is generally true about CPU-side and GPU-side memory?
  5. 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?