Method ForEachEdgeCM

grape/cuda/parallel/parallel_engine.h:1186–1220

  template <typename FRAG_T, typename WORK_SOURCE_T, typename ASSIGN_OP,
            typename EDGE_OP, EdgeDirection ed>
  inline void ForEachEdgeCM(const Stream& stream, const FRAG_T& dev_frag,
                            const WORK_SOURCE_T& ws, ASSIGN_OP assign_op,
                            EDGE_OP op) {
    using vid_t = typename FRAG_T::vid_t;
    using vertex_t = typename FRAG_T::vertex_t;

    int grid_size, block_size;
    size_t size = ws.size();
    if (size == 0) {
      return;
    }

    __calc_prefix_sum__<FRAG_T, WORK_SOURCE_T, ed>(stream, dev_frag, ws);

    auto calc_shmem_size = [] DEV_HOST(int block_size) -> int {
      using vid_t = typename FRAG_T::vid_t;
      using vertex_t = typename FRAG_T::vertex_t;
      using nbr_t = typename FRAG_T::nbr_t;
      using metadata_t = typename std::result_of<ASSIGN_OP&(vertex_t&)>::type;

      return block_size * (sizeof(VertexMetadata<vid_t, metadata_t>) +
                           sizeof(const nbr_t*) + sizeof(vid_t));
    };
    ArrayView<size_t> row_offset_view(
        thrust::raw_pointer_cast(prefix_sum_.data()), size);

    auto lb_wrapper = [=] __device__() {
      LBCM<FRAG_T, WORK_SOURCE_T, ASSIGN_OP, EDGE_OP, ed>(
          dev_frag, ws, row_offset_view, assign_op, op);
    };  // NOLINT

    KernelSizing(grid_size, block_size, ws.size());
    KernelWrapper<<<grid_size, block_size, calc_shmem_size(block_size),
                    stream.cuda_stream()>>>(lb_wrapper);
  }

  template <typename FRAG_T, typename WORK_SOURCE_T, typename ASSIGN_OP,
            typename EDGE_OP, EdgeDirection ed>