Class OrderedSequenceBarrier

include/barrier.h:228–305

template <int SequenceDepth, int SequenceLength>
class OrderedSequenceBarrier {
 public:
  using Barrier = ClusterBarrier;

  struct SharedStorage {
    Barrier barrier_[SequenceDepth][SequenceLength];
  };

  struct Params {
    uint32_t group_id;
    uint32_t group_size;
    int active_warps = 0;
  };

 private:
  // In future this Params object can be replaced easily with a CG object
  Params params_;
  Barrier *barrier_ptr_;
  PipelineState<SequenceDepth> stage_;

  static constexpr int Depth = SequenceDepth;
  static constexpr int Length = SequenceLength;

 public:
  OrderedSequenceBarrier() = delete;
  OrderedSequenceBarrier(const OrderedSequenceBarrier &) = delete;
  OrderedSequenceBarrier(OrderedSequenceBarrier &&) = delete;
  OrderedSequenceBarrier &operator=(const OrderedSequenceBarrier &) = delete;
  OrderedSequenceBarrier &operator=(OrderedSequenceBarrier &&) = delete;
  ~OrderedSequenceBarrier() = default;

  DEVICE OrderedSequenceBarrier(SharedStorage &storage, Params const &params)
      : params_(params),
        barrier_ptr_(&storage.barrier_[0][0]),
        // Group 0 - starts with an opposite phase
        stage_({0, (params.group_id == 0), 0}) {
    int warp_idx = threadIdx.x / WARP_SIZE;
    int lane_predicate = elect_one_sync();

    // Barrier FULL, EMPTY init
    // Init is done only by the one elected thread of the block
    if (warp_idx == params.active_warps && lane_predicate == 1) {
      for (int d = 0; d < Depth; ++d) {
        for (int l = 0; l < Length; ++l) {
          barrier_ptr_[d * Length + l].init(params.group_size);
        }
      }
    }
    fence_barrier_init();
  }

  // Wait on a stage to be unlocked
  DEVICE void wait() {
    get_barrier_for_current_stage(params_.group_id).wait(stage_.phase());
  }

  DEVICE void check_phase(int val) {
    if (threadIdx.x % WARP_GROUP_SIZE == 0) {