Program Listing for File synchronize.hpp

↰ Return to documentation for file (include/ripple/execution/synchronize.hpp)

#ifndef RIPPLE_EXECUTION_SYNCHRONIZE_HPP
#define RIPPLE_EXECUTION_SYNCHRONIZE_HPP

#include <ripple/utility/portability.hpp>
#include <cooperative_groups.h>

namespace ripple {

ripple_all inline auto syncthreads() noexcept -> void {
/*
 * For newer architectures, __syncthreads is called on each thread in a block,
 * so if one thread has returned and __syncthreads is called then there will
 * be a deadlock, which is solved by using the coalesced groups.
 */
#if defined(__CUDA__) && __CUDA_ARCH__ >= 600 && __CUDA_ARCH__ != 700
  //__syncthreads();
  // co-operative groups not working on ampere!
  auto g = cooperative_groups::coalesced_threads();
  g.sync();

/*
 * For older architectures, __syncthreads only needs to succeed on one running
 * thread in the block to avoid deadlock, and the coalesced groupd does not
 * work, so here we default to the old syncthreads.
 */
#elif defined(__CUDA__) && defined(__CUDA_ARCH__)
  __syncthreads();
#endif // __CUDACC__
}

// ripple_all inline auto syncthreads() noexcept -> void {}

} // namespace ripple

#endif // RIPPLE_EXECUTION_SYNCHRONIZE_HPP