Simplified SSSP-Gunrock

test_helpers.hxx

// ----------------------------------------------------------------
// Gunrock -- Fast and Efficient GPU Graph Library
// ----------------------------------------------------------------
// This source code is distributed under the terms of LICENSE.TXT
// in the root directory of this source distribution.
// ----------------------------------------------------------------

/**
 * @file
 * test_helpers.hxx
 *
 * @brief header file for application specific helpers.
 * Note that we may not require all of the includes, most
 * of these can be abstracted away to a single test_util.hxx.
 */

#pragma once

// Utilities and Correctness-Checking
#include <gunrock/util/test_utils.cuh>

// Graph Definitions
#include <gunrock/graphio/graphio.cuh>

#include <gunrock/app/app_base.cuh>
#include <gunrock/app/test_base.cuh>

// #include <gunrock/app/problem_base.cuh>

#include <gunrock/app/enactor_base.cuh>
//#include <gunrock/app/enactor_iteration.cuh>
//#include <gunrock/app/enactor_loop.cuh>

#include <gunrock/app/frontier.cuh>

// Gunrock Operators
#include <gunrock/oprtr/oprtr.cuh>

using namespace gunrock;

/**
 * @brief Displays the SSSP result (i.e., distance from source)
 * @tparam T Type of values to display
 * @tparam SizeT Type of size counters
 * @param[in] preds Search depth from the source for each node.
 * @param[in] num_nodes Number of nodes in the graph.
 */
template <typename T, typename SizeT>
void DisplaySolution(T *array, SizeT length) {
  if (length > 40) length = 40;

  util::PrintMsg("Solution:: [", true, false);
  for (SizeT i = 0; i < length; ++i) {
    util::PrintMsg(std::to_string(i) + ":" + std::to_string(array[i]) + " ",
                   true, false);
  }
  util::PrintMsg("]");
}

template <typename OprtrParametersT, typename SizeT, typename FrontierT>
cudaError_t SetupOprtrParameters(OprtrParametersT &oprtr_parameters,
                                 util::Parameters &parameters,
                                 FrontierT &frontier, SizeT gpu = 0,
                                 int num_gpus = 1) {
  cudaError_t retval = cudaSuccess;
  std::string advance_mode = parameters.Get<std::string>("advance-mode");
  std::string filter_mode = parameters.Get<std::string>("filter-mode");
  int max_grid_size = parameters.Get<int>("max-grid-size");

  cudaStream_t stream = 0;
  cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
  mgpu::ContextPtr context;

  // Device properties
  util::Array1D<SizeT, util::CudaProperties> cuda_props;
  GUARD_CU(cuda_props.Allocate(num_gpus, util::HOST));
  GUARD_CU(util::SetDevice(gpu));

  cuda_props[gpu].Setup(gpu);

  int mgpu_idx;
  GUARD_CU2(cudaGetDevice(&mgpu_idx), "cudaGetDevice failed.");
  context = mgpu::CreateCudaDeviceAttachStream(mgpu_idx, stream);

  util::CudaProperties *cuda_properties = cuda_props + gpu;
  GUARD_CU(oprtr_parameters.Init());

  oprtr_parameters.stream = stream;
  oprtr_parameters.context = context;
  oprtr_parameters.frontier = &frontier;
  oprtr_parameters.cuda_props = cuda_properties;
  oprtr_parameters.advance_mode = advance_mode;
  oprtr_parameters.filter_mode = filter_mode;

  return retval;
}

template <typename FrontierT, typename GraphT>
cudaError_t SetupFrontier(FrontierT &frontier, GraphT &graph,
                          util::Parameters &parameters,
                          util::Location target = util::DEVICE) {
  cudaError_t retval = cudaSuccess;
  typedef typename GraphT::VertexT VertexT;

  std::vector<double> queue_factors =
      parameters.Get<std::vector<double>>("queue-factor");

  std::string frontier_name = "app::frontier[0]";
  int num_queues = 3;

  GUARD_CU(frontier.Init(num_queues, NULL, frontier_name, target));
  GUARD_CU(frontier.Allocate(graph.nodes, graph.edges, queue_factors));

  int peer_ = 0;
  frontier.queue_length = (peer_ == 0) ? 1 : 0;
  printf("frontier.queue_length: %lu\n", frontier.queue_length);

  return retval;
}

test_sssp.cu

// ----------------------------------------------------------------
// Gunrock -- Fast and Efficient GPU Graph Library
// ----------------------------------------------------------------
// This source code is distributed under the terms of LICENSE.TXT
// in the root directory of this source distribution.
// ----------------------------------------------------------------

/**
 * @file
 * test_sssp.cu
 *
 * @brief Simple test driver program for single source shortest path.
 */

#include "test_helpers.hxx"

using namespace gunrock;

cudaError_t UseParameters(util::Parameters &parameters) {
  cudaError_t retval = cudaSuccess;

  app::UseParameters_enactor(parameters);

  GUARD_CU(parameters.Use<std::string>(
      "src",
      util::REQUIRED_ARGUMENT | util::MULTI_VALUE | util::OPTIONAL_PARAMETER,
      "0",
      "<Vertex-ID|random|largestdegree> The source vertices\n"
      "\tIf random, randomly select non-zero degree vertices;\n"
      "\tIf largestdegree, select vertices with largest degrees",
      __FILE__, __LINE__));

  GUARD_CU(parameters.Use<int>(
      "src-seed",
      util::REQUIRED_ARGUMENT | util::SINGLE_VALUE | util::OPTIONAL_PARAMETER,
      util::PreDefinedValues<int>::InvalidValue,
      "seed to generate random sources", __FILE__, __LINE__));

  return retval;
}

template <typename GraphT, typename FrontierT, typename OprtrParametersT>
cudaError_t SSSP(GraphT &graph, util::Parameters &parameters,
                 FrontierT &frontier, OprtrParametersT &oprtr_parameters,
                 util::Location target = util::DEVICE) {
  cudaError_t retval = cudaSuccess;

  typedef typename GraphT::VertexT VertexT;
  typedef typename GraphT::VertexT LabelT;
  typedef typename GraphT::SizeT SizeT;
  typedef typename GraphT::ValueT ValueT;
  typedef typename GraphT::CsrT CsrT;

  bool quiet = parameters.Get<bool>("quiet");
  std::vector<VertexT> srcs = parameters.Get<std::vector<VertexT>>("srcs");

  util::CpuTimer cpu_timer;
  auto &stream = oprtr_parameters.stream;
  auto nodes = graph.nodes;
  auto edges = graph.edges;
  int num_srcs = srcs.size();
  VertexT src = srcs[0];

  util::Array1D<SizeT, ValueT> distances;  // source distance
  util::Array1D<SizeT, LabelT>
      labels;  // labels to mark latest iteration the vertex been visited
  auto &weights = graph.CsrT::edge_values;

  GUARD_CU(graph.Move(util::HOST, target, stream));

  // Allocate memory for distances and labels on the GPU (or target)
  GUARD_CU(distances.Allocate(nodes, util::HOST | target));
  GUARD_CU(labels.Allocate(nodes, target));

  // Reset data
  GUARD_CU(distances.ForEach(
      [] __host__ __device__(ValueT & distance) {
        distance = util::PreDefinedValues<ValueT>::MaxValue;
      },
      nodes, target, stream));

  GUARD_CU(labels.ForEach(
      [] __host__ __device__(LabelT & label) {
        label = util::PreDefinedValues<LabelT>::InvalidValue;
      },
      nodes, target, stream));

  GUARD_CU(distances.ForAll(
      [src] __host__ __device__(ValueT * distances_, const SizeT &pos) {
        distances_[src] = 0;
      },
      1, target));

  GUARD_CU(frontier.V_Q()->ForEach(
      [src] __host__ __device__(VertexT & v) { v = src; }, 1, target, stream));

  cpu_timer.Start();
  SizeT iteration = 0;
  while (frontier.queue_length > 0) {
    // The advance operation
    auto advance_op = [distances, weights] __host__ __device__(
                          const VertexT &src, VertexT &dest,
                          const SizeT &edge_id, const VertexT &input_item,
                          const SizeT &input_pos, SizeT &output_pos) -> bool {
      ValueT src_distance = Load<cub::LOAD_CG>(distances + src);
      ValueT edge_weight = Load<cub::LOAD_CS>(weights + edge_id);
      ValueT new_distance = src_distance + edge_weight;

      // Check if the destination node has been claimed as someone's child
      ValueT old_distance = atomicMin(distances + dest, new_distance);

      if (new_distance < old_distance) {
        return true;
      }
      return false;
    };

    // The filter operation
    auto filter_op = [labels, iteration] __host__ __device__(
                         const VertexT &src, VertexT &dest,
                         const SizeT &edge_id, const VertexT &input_item,
                         const SizeT &input_pos, SizeT &output_pos) -> bool {
      if (!util::isValid(dest)) return false;
      if (labels[dest] == iteration) return false;
      labels[dest] = iteration;
      return true;
    };

    oprtr_parameters.label = iteration + 1;
    // Call the advance operator, using the advance operation
    GUARD_CU(oprtr::Advance<oprtr::OprtrType_V2V>(
        graph.csr(), frontier.V_Q(), frontier.Next_V_Q(), oprtr_parameters,
        advance_op, filter_op));

    if (oprtr_parameters.advance_mode != "LB_CULL" &&
        oprtr_parameters.advance_mode != "LB_LIGHT_CULL") {
      frontier.queue_reset = false;
      // Call the filter operator, using the filter operation
      GUARD_CU(oprtr::Filter<oprtr::OprtrType_V2V>(
          graph.csr(), frontier.V_Q(), frontier.Next_V_Q(), oprtr_parameters,
          filter_op));
    }

    // Get back the resulted frontier length
    GUARD_CU(frontier.work_progress.GetQueueLength(
        frontier.queue_index, frontier.queue_length, false,
        oprtr_parameters.stream, true));
    iteration++;
  }
  cpu_timer.Stop();
  float elapsed_time = cpu_timer.ElapsedMillis();

  util::PrintMsg("Elapsed Time: " + std::to_string(elapsed_time) + " ms", !quiet);

  GUARD_CU(distances.Move(util::DEVICE, util::HOST));
  DisplaySolution(distances.GetPointer(util::HOST), nodes);

  return retval;
}

/******************************************************************************
 * Main
 ******************************************************************************/

/**
 * @brief Enclosure to the main function
 */
struct main_struct {
  /**
   * @brief the actual main function, after type switching
   * @tparam VertexT    Type of vertex identifier
   * @tparam SizeT      Type of graph size, i.e. type of edge identifier
   * @tparam ValueT     Type of edge values
   * @param  parameters Command line parameters
   * @param  v,s,val    Place holders for type deduction
   * \return cudaError_t error message(s), if any
   */
  template <typename VertexT,  // Use int as the vertex identifier
            typename SizeT,    // Use int as the graph size type
            typename ValueT>   // Use int as the value type
  cudaError_t operator()(util::Parameters &parameters, VertexT v, SizeT s,
                         ValueT val) {
    typedef typename app::TestGraph<VertexT, SizeT, ValueT,
                                    graph::HAS_EDGE_VALUES | graph::HAS_CSR>
        GraphT;
    typedef typename GraphT::CsrT CsrT;
    typedef typename GraphT::VertexT LabelT;
    typedef app::Frontier<VertexT, SizeT> FrontierT;
    typedef oprtr::OprtrParameters<GraphT, FrontierT, LabelT> OprtrParametersT;

    cudaError_t retval = cudaSuccess;

    std::vector<int> gpu_idx = parameters.Get<std::vector<int>>("device");
    bool quiet = parameters.Get<bool>("quiet");

    int num_gpus = gpu_idx.size();
    if (num_gpus > 1 || num_gpus < 1) num_gpus = 1;

    util::CpuTimer cpu_timer;
    GraphT graph;

    cpu_timer.Start();
    GUARD_CU(graphio::LoadGraph(parameters, graph));
    cpu_timer.Stop();
    float load_time = cpu_timer.ElapsedMillis();
    parameters.Set("load-time", load_time);

    GUARD_CU(app::Set_Srcs(parameters, graph));

    FrontierT frontier;
    SetupFrontier(frontier, graph, parameters);

    OprtrParametersT oprtr_parameters;
    SetupOprtrParameters(oprtr_parameters, parameters, frontier, gpu_idx[0],
                         num_gpus);


    util::PrintMsg("[ SSSP Application Statistics ]", !quiet);
    util::PrintMsg("-------------------------------", !quiet);

    cpu_timer.Start();
    SSSP(graph, parameters, frontier, oprtr_parameters);
    cpu_timer.Stop();

    util::PrintMsg("Total Time: "
	+ std::to_string(cpu_timer.ElapsedMillis()) + " ms", !quiet);
    util::PrintMsg("Load Time: "
	+ std::to_string(load_time) + " ms", !quiet);

    return retval;
  }
};

int main(int argc, char **argv) {
  cudaError_t retval = cudaSuccess;
  util::Parameters parameters("test sssp");
  GUARD_CU(graphio::UseParameters(parameters));
  GUARD_CU(UseParameters(parameters));
  GUARD_CU(app::UseParameters_test(parameters));
  GUARD_CU(parameters.Parse_CommandLine(argc, argv));
  if (parameters.Get<bool>("help")) {
    parameters.Print_Help();
    return cudaSuccess;
  }
  GUARD_CU(parameters.Check_Required());

  return app::Switch_Types<app::VERTEXT_U32B | app::VERTEXT_U64B |
                           app::SIZET_U32B | app::SIZET_U64B |
                           app::VALUET_U32B | app::DIRECTED | app::UNDIRECTED>(
      parameters, main_struct());
}

// Leave this at the end of the file
// Local Variables:
// mode:c++
// c-file-style: "NVIDIA"
// End: