Can I use __syncthreads() after having dropped threads?

by

The answer to the short question is “No”. Warp level branch divergence around a __syncthreads() instruction will cause a deadlock and result in a kernel hang. Your code example is not guaranteed to be safe or correct. The correct way to implement the code would be like this:

__global__ void kernel(...)

    if (tidx < N) {
        // Code stanza #1
    }

    __syncthreads();


    if (tidx < N) {
        // Code stanza #2
    }

    // etc
}

so that the __syncthreads() instructions are executed unconditionally.

EDIT: Just to add a bit of additional information which confirms this assertion, __syncthreads() calls get compiled into the PTX bar.sync instruction on all architectures. The PTX2.0 guide (p133) documents bar.sync and includes the following warning:

Barriers are executed on a per-warp basis as if all the threads in a
warp are active. Thus, if any thread in a warp executes a bar
instruction, it is as if all the threads in the warp have executed the
bar instruction. All threads in the warp are stalled until the barrier
completes, and the arrival count for the barrier is incremented by the
warp size (not the number of active threads in the warp). In
conditionally executed code, a bar instruction should only be used if
it is known that all threads evaluate the condition identically (the
warp does not diverge). Since barriers are executed on a per-warp
basis, the optional thread count must be a multiple of the warp size.

So despite any assertions to the contrary, it is not safe to have conditional branching around a __syncthreads() call unless you can be 100% certain that every thread in any given warp follows the same code path and no warp divergence can occur.

More Related Contents:

  1. When should one use a spinlock instead of mutex?
  2. Dependent loads reordering in CPU
  3. Symbolic links and synced folders in Vagrant
  4. Asynchronous Process inside a javascript for loop [duplicate]
  5. How do I choose grid and block dimensions for CUDA kernels?
  6. How to synchronize a static variable among threads running different instances of a class in Java?
  7. Using Java with Nvidia GPUs (CUDA)
  8. Java double checked locking
  9. Synchronizing on an Integer value [duplicate]
  10. Thrust inside user written kernels
  11. CUDA, Using 2D and 3D Arrays
  12. sending 3d array to CUDA kernel
  13. Passing pointers between C and Java through JNI
  14. Why does cudaMalloc() use pointer to pointer?
  15. How do CUDA blocks/warps/threads map onto CUDA cores?
  16. Understanding CUDA grid dimensions, block dimensions and threads organization (simple explanation) [closed]
  17. Sharing a variable between multiple different threads
  18. Using string as a lock to do thread synchronization
  19. How to wait until a predicate condition becomes true in JavaScript?
  20. fs.writeFile in a promise, asynchronous-synchronous stuff
  21. CUDA apps time out & fail after several seconds – how to work around this?
  22. Why is NVIDIA Pascal GPUs slow on running CUDA Kernels when using cudaMallocManaged
  23. How to acquire a lock by a key
  24. Sync JS time between multiple devices
  25. std::mutex performance compared to win32 CRITICAL_SECTION
  26. How to create a CUDA context?
  27. Why does my CUDA kernel crash (unspecified launch failure) with a different dataset size?
  28. ReleaseSemaphore does not release the semaphore
  29. printf() in my CUDA kernel doesn’t result produce any output
  30. “invalid configuration argument ” error for the call of CUDA kernel?

Leave a Comment