GPU Programming
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)?
if (threadIdx.x % 2 == 0) A(); else B(); // (1)
if ((threadIdx.x / 32) % 2 == 0) A(); else B(); // (2)- A(1) is more efficient because its branches agree
- B(1) diverges inside every warp, while (2) is uniform per warp, so (2) is more efficient
- CBoth are exactly equally efficient
Show answer and explanation
Answer: B. (1) diverges inside every warp, while (2) is uniform per warp, so (2) is more efficient
In (1), even and odd threads split, so every warp contains both A and B threads and every warp executes A and B in turn. In (2), the split is by the thread index divided by 32, so all 32 threads of a warp take the same side, and each warp executes only A or only B. The same work avoids divergence when the condition is uniform per warp.
- A. In (1), adjacent threads have different conditions.
- B. Correct. The trick is to make the condition uniform per warp.
- C. They differ greatly in whether divergence occurs.