GPU Programming

In CUDA, __syncthreads() is placed inside "if (threadIdx.x < 16)" so that only some threads reach it. What can happen?

  1. ASynchronization works correctly among just the 16 threads
  2. BThe other threads never arrive, so the block may hang or results may be undefined
  3. C__syncthreads() is ignored automatically, so nothing happens
Show answer and explanation

Answer: B. The other threads never arrive, so the block may hang or results may be undefined

__syncthreads() waits for every thread in the block, so placing it inside a branch only some threads take means some threads never arrive. The behavior is undefined by the specification; in practice it may hang or appear to work while producing wrong results. Synchronization should be placed outside branches, where every thread is guaranteed to pass.

  • A. It waits for the whole block.
  • B. Correct. It must be placed where every thread in the block passes.
  • C. It is not ignored.
Practice this topic Back to topic

More GPU Programming questions

  1. Which describes a bank conflict in CUDA shared memory?
  2. 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?
  3. When multiplying large matrices on a GPU, why are the matrices split into small tiles loaded into shared memory?
  4. What can happen if a GPU kernel uses too many registers per thread?
  5. Which describes occupancy on a GPU?