tags: cuda programming, programming, deep learning
All threads that are running in a block, are all loaded in a SM simultaneously. The threads can be hanged to wait for each other using the __syncthreads() call when they are in SAME BLOCK .
preventing threads in different blocks from communicating with barrier synchronization allows the different blocks do not have a certain execution order and them not relying on each other, which leads to transparent scalability i.e. running the same application on different devices with different parallelization ability.
The awkward if-then
The book wrote a weird if-then condition
if (threadsIdx.x % 2 == 0) {
__syncthreads();
} else {
__sycthreads();
}Even though this code block looks relatively innocent, according to the book, we can not do that. Every __syncthread() call sets up a synchronization barrier in that specific location, either all the threads in the block branch to the if-then, or all of them branch to the else. They can not be mixed because they all must hit the same specific synchronization point. TODO: understand better and fact check