Is the statement "All the threads that run on a given SM may cooperate" true or false?

Question

In the Quiz: Thread and Blocks of "Intro to Parallel Programming" on Udacity, it asks whether the statement "All the threads that run on a given SM may cooperate to solve a problem" is true or false. The answer is false however I think it should be true because:

• It use "may" rather than "must".
• Some paper are talking about inter-block GPU communication.

Do I have any misunderstanding of it?

Answer 1

It is not true in all cases.

In the typical CUDA sense of cooperating threads, cooperating threads must belong to the same block (so as to be able to use cooperative mechanisms like __syncthreads() and __shared__ memory).

An SM, however, can contain multiple blocks executing simultaneously.

Threads from separate blocks, even if resident on the same SM, cannot cooperate, in the typical CUDA sense of the words "cooperating threads".

Answer 2

The CUDA official material[1] mentioned what Robert Crovella said in page 8:

• Threads in different blocks cannot cooperate.

[1] https://www.nvidia.com/content/cudazone/download/Getting_Started_w_CUDA_Training_NVISION08.pdf