Skip to content

Commit

Permalink
fixed the interface mess. managed to use only one pass of exclusive s…
Browse files Browse the repository at this point in the history 
…can in advance.
  • Loading branch information
Yangzihao Wang committed Feb 5, 2015
1 parent f63ac6f commit b588f89
Show file tree
Hide file tree
Showing 6 changed files with 33 additions and 29 deletions.
16 changes: 10 additions & 6 deletions gunrock/app/bfs/bfs_problem.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -32,16 +32,20 @@ namespace bfs {
* @tparam _USE_DOUBLE_BUFFER Boolean type parameter which defines whether to use double buffer.
*/
template <
typename VertexId,
typename SizeT,
typename Value,
typename _VertexId,
typename _SizeT,
typename _Value,
bool _MARK_PREDECESSORS,
bool _ENABLE_IDEMPOTENCE,
bool _USE_DOUBLE_BUFFER>
struct BFSProblem : ProblemBase<VertexId, SizeT,
struct BFSProblem : ProblemBase<_VertexId, _SizeT,
_USE_DOUBLE_BUFFER>
{

typedef _VertexId VertexId;
typedef _SizeT SizeT;
typedef _Value Value;

static const bool MARK_PREDECESSORS = _MARK_PREDECESSORS;
static const bool ENABLE_IDEMPOTENCE = _ENABLE_IDEMPOTENCE;

Expand Down Expand Up @@ -191,7 +195,7 @@ struct BFSProblem : ProblemBase<VertexId, SizeT,
edges = graph.edges;
VertexId *h_row_offsets = graph.row_offsets;
VertexId *h_column_indices = graph.column_indices;
ProblemBase<VertexId, SizeT, _USE_DOUBLE_BUFFER>::Init(stream_from_host,
ProblemBase<_VertexId, _SizeT, _USE_DOUBLE_BUFFER>::Init(stream_from_host,
nodes,
edges,
h_row_offsets,
Expand Down Expand Up @@ -270,7 +274,7 @@ struct BFSProblem : ProblemBase<VertexId, SizeT,
FrontierType frontier_type, // The frontier type (i.e., edge/vertex/mixed)
double queue_sizing) // Size scaling factor for work queue allocation (e.g., 1.0 creates n-element and m-element vertex and edge frontiers, respectively). 0.0 is unspecified.
{
typedef ProblemBase<VertexId, SizeT, _USE_DOUBLE_BUFFER> BaseProblem;
typedef ProblemBase<_VertexId, _SizeT, _USE_DOUBLE_BUFFER> BaseProblem;
//load ProblemBase Reset
BaseProblem::Reset(frontier_type, queue_sizing);

Expand Down
26 changes: 13 additions & 13 deletions gunrock/oprtr/advance/kernel.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -93,8 +93,8 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
CudaContext &context,
TYPE ADVANCE_TYPE,
bool inverse_graph = false,
OP REDUCE_OP = KernelPolicy::OP::NONE,
REDUCE_TYPE R_TYPE = KernelPolicy::REDUCE_TYPE::NONE,
OP REDUCE_OP = gunrock::oprtr::advance::NONE,
REDUCE_TYPE R_TYPE = gunrock::oprtr::advance::EMPTY,
typename KernelPolicy::Value *d_value_to_reduce = NULL,
typename KernelPolicy::Value *d_reduce_frontier = NULL,
typename KernelPolicy::Value *d_reduced_value = NULL)
Expand Down Expand Up @@ -146,16 +146,16 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
d_row_indices,
d_in_key_queue,
partitioned_scanned_edges,
frontier_attribute.queue_length,
frontier_attribute.queue_length+1,
max_in,
max_out,
ADVANCE_TYPE);

Scan<mgpu::MgpuScanTypeInc>((int*)partitioned_scanned_edges, frontier_attribute.queue_length, (int)0, mgpu::plus<int>(),
Scan<mgpu::MgpuScanTypeExc>((int*)partitioned_scanned_edges, frontier_attribute.queue_length+1, (int)0, mgpu::plus<int>(),
(int*)0, (int*)0, (int*)partitioned_scanned_edges, context);

SizeT *temp = new SizeT[1];
cudaMemcpy(temp,partitioned_scanned_edges+frontier_attribute.queue_length-1, sizeof(SizeT), cudaMemcpyDeviceToHost);
cudaMemcpy(temp,partitioned_scanned_edges+frontier_attribute.queue_length, sizeof(SizeT), cudaMemcpyDeviceToHost);
SizeT output_queue_len = temp[0];
//printf("input queue:%d, output_queue:%d\n", frontier_attribute.queue_length, output_queue_len);

Expand All @@ -169,7 +169,7 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
d_column_offsets,
d_row_indices,
(VertexId*)NULL,
partitioned_scanned_edges,
&partitioned_scanned_edges[1],
d_done,
d_in_key_queue,
backward_frontier_map_in,
Expand All @@ -193,7 +193,7 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
d_column_offsets,
d_row_indices,
(VertexId*)NULL,
partitioned_scanned_edges,
&partitioned_scanned_edges[1],
d_done,
d_in_key_queue,
backward_frontier_map_out,
Expand Down Expand Up @@ -266,16 +266,16 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
d_column_indices,
d_in_key_queue,
partitioned_scanned_edges,
frontier_attribute.queue_length,
frontier_attribute.queue_length+1,
max_in,
max_out,
ADVANCE_TYPE);

Scan<mgpu::MgpuScanTypeInc>((int*)partitioned_scanned_edges, frontier_attribute.queue_length, (int)0, mgpu::plus<int>(),
Scan<mgpu::MgpuScanTypeExc>((int*)partitioned_scanned_edges, frontier_attribute.queue_length+1, (int)0, mgpu::plus<int>(),
(int*)0, (int*)0, (int*)partitioned_scanned_edges, context);

SizeT *temp = new SizeT[1];
cudaMemcpy(temp,partitioned_scanned_edges+frontier_attribute.queue_length-1, sizeof(SizeT), cudaMemcpyDeviceToHost);
cudaMemcpy(temp,partitioned_scanned_edges+frontier_attribute.queue_length, sizeof(SizeT), cudaMemcpyDeviceToHost);
SizeT output_queue_len = temp[0];

if (output_queue_len < LBPOLICY::LIGHT_EDGE_THRESHOLD)
Expand All @@ -288,7 +288,7 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
d_row_offsets,
d_column_indices,
d_row_indices,
partitioned_scanned_edges,
&partitioned_scanned_edges[1],
d_done,
d_in_key_queue,
d_out_key_queue,
Expand Down Expand Up @@ -321,7 +321,7 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
SortedSearch<MgpuBoundsLower>(
enactor_stats.d_node_locks,
KernelPolicy::LOAD_BALANCED::BLOCKS,
partitioned_scanned_edges,
&partitioned_scanned_edges[1],
frontier_attribute.queue_length,
enactor_stats.d_node_locks_out,
context);
Expand All @@ -334,7 +334,7 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
d_row_offsets,
d_column_indices,
d_row_indices,
partitioned_scanned_edges,
&partitioned_scanned_edges[1],
enactor_stats.d_node_locks_out,
KernelPolicy::LOAD_BALANCED::BLOCKS,
d_done,
Expand Down
1 change: 1 addition & 0 deletions gunrock/oprtr/advance/kernel_policy.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -132,6 +132,7 @@ struct KernelPolicy {
typedef _ProblemData ProblemData;
typedef typename ProblemData::VertexId VertexId;
typedef typename ProblemData::SizeT SizeT;
typedef typename ProblemData::Value Value;

static const MODE ADVANCE_MODE = _ADVANCE_MODE;
static const int CTA_OCCUPANCY = _MIN_CTA_OCCUPANCY;
Expand Down
11 changes: 6 additions & 5 deletions gunrock/oprtr/edge_map_partitioned/kernel.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -196,14 +196,15 @@ struct Dispatch<KernelPolicy, ProblemData, Functor, true>
int bid = blockIdx.x;

int my_id = bid*blockDim.x + tid;
if (my_id >= num_elements || my_id >= max_edge)
if (my_id > num_elements || my_id >= max_edge)
return;
VertexId v_id = d_queue[my_id];
if (v_id == -1) {
d_scanned_edges[my_id] = 0;
return;
}
SizeT num_edges = GetNeighborListLength(d_row_offsets, d_column_indices, v_id, max_vertex, max_edge, ADVANCE_TYPE);
// add a zero length neighbor list to the end (this for getting both exclusive and inclusive scan in one array)
SizeT num_edges = (my_id == num_elements) ? 0 : GetNeighborListLength(d_row_offsets, d_column_indices, v_id, max_vertex, max_edge, ADVANCE_TYPE);
d_scanned_edges[my_id] = num_edges;
}

Expand Down Expand Up @@ -591,7 +592,7 @@ struct Dispatch<KernelPolicy, ProblemData, Functor, true>
util::CtaWorkProgress &work_progress,
util::KernelRuntimeStats &kernel_stats,
gunrock::oprtr::advance::TYPE &ADVANCE_TYPE,
bool inverse_graph,
bool &inverse_graph,
gunrock::oprtr::advance::REDUCE_TYPE R_TYPE,
Value *&d_value_to_reduce,
Value *&d_reduce_frontier)
Expand Down Expand Up @@ -874,8 +875,8 @@ void RelaxPartitionedEdges2(
gunrock::oprtr::advance::TYPE ADVANCE_TYPE = gunrock::oprtr::advance::V2V,
bool inverse_graph = false,
gunrock::oprtr::advance::REDUCE_TYPE R_TYPE = gunrock::oprtr::advance::EMPTY,
typename KernelPolicy::Value d_value_to_reduce = NULL,
typename KernelPolicy::Value d_reduce_frontier = NULL)
typename KernelPolicy::Value *d_value_to_reduce = NULL,
typename KernelPolicy::Value *d_reduce_frontier = NULL)
{
Dispatch<KernelPolicy, ProblemData, Functor>::RelaxPartitionedEdges2(
queue_reset,
Expand Down
1 change: 1 addition & 0 deletions gunrock/oprtr/edge_map_partitioned/kernel_policy.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -69,6 +69,7 @@ struct KernelPolicy
typedef _ProblemData ProblemData;
typedef typename ProblemData::VertexId VertexId;
typedef typename ProblemData::SizeT SizeT;
typedef typename ProblemData::Value Value;

enum {

Expand Down
7 changes: 2 additions & 5 deletions tests/bfs/test_bfs.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -659,11 +659,8 @@ int main( int argc, char** argv)
csr.PrintHistogram();

// Run tests
//RunTests(csr, args, *context);
int temp;
int src = csr.GetNodeWithHighestDegree(temp);
printf("highest degree:%d\n", temp);
} else {
RunTests(csr, args, *context);
} else {

// Unknown graph type
fprintf(stderr, "Unspecified graph type\n");
Expand Down

0 comments on commit b588f89

Please sign in to comment.