Skip to content

Commit

Permalink
Merge pull request #7 from tahsinreza/dijkstra-rename
Browse files Browse the repository at this point in the history 
Renamed Dijkstra instances to SSSP and updated the new SSSP unit-test file accordingly. Deleted old Dijkstra algorithm and unit-test files.
  • Loading branch information
@ahgharaibeh
ahgharaibeh committed Aug 27, 2014
2 parents 2e01748 + 8e80b45 commit 4b10dc9
Show file tree
Hide file tree
Showing 8 changed files with 43 additions and 321 deletions.
8 changes: 4 additions & 4 deletions src/alg/totem_alg.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -152,10 +152,10 @@ error_t clustering_coefficient_sorted_neighbours_gpu(const graph_t* graph,
* @param[out] shortest_distances the length of the computed shortest paths
* @return generic success or failure
*/
error_t dijkstra_cpu(const graph_t* graph, vid_t src_id, weight_t* distance);
error_t dijkstra_gpu(const graph_t* graph, vid_t src_id, weight_t* distance);
error_t dijkstra_vwarp_gpu(const graph_t* graph, vid_t src_id,
weight_t* distance);
error_t sssp_cpu(const graph_t* graph, vid_t src_id, weight_t* distance);
error_t sssp_gpu(const graph_t* graph, vid_t src_id, weight_t* distance);
error_t sssp_vwarp_gpu(const graph_t* graph, vid_t src_id,
weight_t* distance);
error_t sssp_hybrid(vid_t src_id, weight_t* distance);

/**
Expand Down
8 changes: 4 additions & 4 deletions src/alg/totem_apsp.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,8 +13,8 @@
#include "totem_alg.h"

// Externed function declarations for dijkstra kernel functions
__global__ void dijkstra_kernel(graph_t, bool*, weight_t*, weight_t*);
__global__ void dijkstra_final_kernel(graph_t, bool*, weight_t*, weight_t*,
__global__ void sssp_kernel(graph_t, bool*, weight_t*, weight_t*);
__global__ void sssp_final_kernel(graph_t, bool*, weight_t*, weight_t*,
bool* has_true);


Expand Down Expand Up @@ -222,12 +222,12 @@ error_t apsp_gpu(graph_t* graph, weight_t** path_ret) {
new_distances_d), err_free_all);
bool has_true = true;
while (has_true) {
dijkstra_kernel<<<block_count, threads_per_block>>>
sssp_kernel<<<block_count, threads_per_block>>>
(*graph_d, changed_d, distances_d, new_distances_d);
CHK_CU_SUCCESS(cudaMemset(changed_d, false, graph->vertex_count *
sizeof(bool)), err_free_all);
CHK_CU_SUCCESS(cudaMemset(has_true_d, false, sizeof(bool)), err_free_all);
dijkstra_final_kernel<<<block_count, threads_per_block>>>
sssp_final_kernel<<<block_count, threads_per_block>>>
(*graph_d, changed_d, distances_d, new_distances_d, has_true_d);
CHK_CU_SUCCESS(cudaMemcpy(&has_true, has_true_d, sizeof(bool),
cudaMemcpyDeviceToHost), err_free_all);
Expand Down
22 changes: 11 additions & 11 deletions src/alg/totem_dijkstra.cu → src/alg/totem_sssp.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -155,7 +155,7 @@ error_t finalize_gpu(graph_t* graph_d, weight_t* distances_d, bool* changed_d,
* @param[out] new_distances an array with distances updated in this round
*/
__global__
void dijkstra_kernel(graph_t graph, bool* to_update, weight_t* distances,
void sssp_kernel(graph_t graph, bool* to_update, weight_t* distances,
weight_t* new_distances) {
const vid_t vertex_id = THREAD_GLOBAL_INDEX;
if ((vertex_id >= graph.vertex_count) || !to_update[vertex_id]) {
Expand Down Expand Up @@ -198,7 +198,7 @@ void vwarp_process_neighbors(vid_t warp_offset, vid_t neighbor_count,
* technique.
*/
__global__
void vwarp_dijkstra_kernel(graph_t graph, bool* to_update, weight_t* distances,
void vwarp_sssp_kernel(graph_t graph, bool* to_update, weight_t* distances,
weight_t* new_distances, uint32_t thread_count) {

if (THREAD_GLOBAL_INDEX >= thread_count) return;
Expand Down Expand Up @@ -243,7 +243,7 @@ void vwarp_dijkstra_kernel(graph_t graph, bool* to_update, weight_t* distances,
* @param[out] new_distances an array with distances updated in this round
*/
__global__
void dijkstra_final_kernel(graph_t graph, bool* to_update, weight_t* distances,
void sssp_final_kernel(graph_t graph, bool* to_update, weight_t* distances,
weight_t* new_distances, bool* has_true) {
const vid_t vertex_id = THREAD_GLOBAL_INDEX;
if (vertex_id >= graph.vertex_count) {
Expand All @@ -257,7 +257,7 @@ void dijkstra_final_kernel(graph_t graph, bool* to_update, weight_t* distances,
new_distances[vertex_id] = distances[vertex_id];
}

error_t dijkstra_gpu(const graph_t* graph, vid_t source_id,
error_t sssp_gpu(const graph_t* graph, vid_t source_id,
weight_t* shortest_distances) {
// Check for special cases
bool finished = false;
Expand All @@ -280,12 +280,12 @@ error_t dijkstra_gpu(const graph_t* graph, vid_t source_id,
KERNEL_CONFIGURE(graph->vertex_count, block_count, threads_per_block);
bool has_true = true;
while (has_true) {
dijkstra_kernel<<<block_count, threads_per_block>>>
sssp_kernel<<<block_count, threads_per_block>>>
(*graph_d, changed_d, distances_d, new_distances_d);
CHK_CU_SUCCESS(cudaMemset(changed_d, false, graph->vertex_count *
sizeof(bool)), err_free_all);
CHK_CU_SUCCESS(cudaMemset(has_true_d, false, sizeof(bool)), err_free_all);
dijkstra_final_kernel<<<block_count, threads_per_block>>>
sssp_final_kernel<<<block_count, threads_per_block>>>
(*graph_d, changed_d, distances_d, new_distances_d, has_true_d);
CHK_CU_SUCCESS(cudaMemcpy(&has_true, has_true_d, sizeof(bool),
cudaMemcpyDeviceToHost), err_free_all);
Expand All @@ -309,7 +309,7 @@ error_t dijkstra_gpu(const graph_t* graph, vid_t source_id,
return FAILURE;
}

error_t dijkstra_vwarp_gpu(const graph_t* graph, vid_t source_id,
error_t sssp_vwarp_gpu(const graph_t* graph, vid_t source_id,
weight_t* shortest_distances) {

// Check for special cases
Expand All @@ -334,18 +334,18 @@ error_t dijkstra_vwarp_gpu(const graph_t* graph, vid_t source_id,
bool has_true = true;
dim3 block_count, threads_per_block;
KERNEL_CONFIGURE(thread_count, block_count, threads_per_block);
cudaFuncSetCacheConfig(vwarp_dijkstra_kernel, cudaFuncCachePreferShared);
cudaFuncSetCacheConfig(vwarp_sssp_kernel, cudaFuncCachePreferShared);
dim3 block_count_final, threads_per_block_final;
KERNEL_CONFIGURE(graph->vertex_count, block_count_final,
threads_per_block_final);
while (has_true) {
vwarp_dijkstra_kernel<<<block_count, threads_per_block>>>
vwarp_sssp_kernel<<<block_count, threads_per_block>>>
(*graph_d, changed_d, distances_d, new_distances_d, thread_count);
CHK_CU_SUCCESS(cudaMemset(changed_d, false,
vwarp_default_state_length(graph->vertex_count) *
sizeof(bool)), err_free_all);
CHK_CU_SUCCESS(cudaMemset(has_true_d, false, sizeof(bool)), err_free_all);
dijkstra_final_kernel<<<block_count_final, threads_per_block_final>>>
sssp_final_kernel<<<block_count_final, threads_per_block_final>>>
(*graph_d, changed_d, distances_d, new_distances_d, has_true_d);
CHK_CU_SUCCESS(cudaMemcpy(&has_true, has_true_d, sizeof(bool),
cudaMemcpyDeviceToHost), err_free_all);
Expand All @@ -368,7 +368,7 @@ error_t dijkstra_vwarp_gpu(const graph_t* graph, vid_t source_id,
return FAILURE;
}

__host__ error_t dijkstra_cpu(const graph_t* graph, vid_t source_id,
__host__ error_t sssp_cpu(const graph_t* graph, vid_t source_id,
weight_t* shortest_distances) {
// Check for special cases
bool finished = false;
Expand Down
2 changes: 1 addition & 1 deletion src/benchmark/totem_benchmark.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@
typedef enum {
BENCHMARK_BFS = 0,
BENCHMARK_PAGERANK,
BENCHMARK_DIJKSTRA,
BENCHMARK_SSSP,
BENCHMARK_BETWEENNESS,
BENCHMARK_GRAPH500,
BENCHMARK_CLUSTERING_COEFFICIENT,
Expand Down
14 changes: 7 additions & 7 deletions src/benchmark/totem_benchmark_binary.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@
// Defines attributes of the algorithms available for benchmarking
PRIVATE void benchmark_bfs(graph_t*, void*, totem_attr_t*);
PRIVATE void benchmark_pagerank(graph_t*, void*, totem_attr_t*);
PRIVATE void benchmark_dijkstra(graph_t*, void*, totem_attr_t*);
PRIVATE void benchmark_sssp(graph_t*, void*, totem_attr_t*);
PRIVATE void benchmark_betweenness(graph_t*, void*, totem_attr_t*);
PRIVATE void benchmark_graph500(graph_t* graph, void* tree, totem_attr_t* attr);
PRIVATE void benchmark_clustering_coefficient(graph_t* graph, void*,
Expand All @@ -39,8 +39,8 @@ const benchmark_attr_t BENCHMARKS[] = {
NULL
},
{
benchmark_dijkstra,
"DIJKSTRA",
benchmark_sssp,
"SSSP",
sizeof(weight_t),
true,
sizeof(weight_t) * BITS_PER_BYTE,
Expand Down Expand Up @@ -101,7 +101,7 @@ PRIVATE uint64_t get_traversed_edges(graph_t* graph, void* benchmark_output) {
}
break;
}
case BENCHMARK_DIJKSTRA: {
case BENCHMARK_SSSP: {
OMP(omp parallel for reduction(+ : trv_edges))
for (vid_t vid = 0; vid < graph->vertex_count; vid++) {
weight_t* distance = reinterpret_cast<weight_t*>(benchmark_output);
Expand Down Expand Up @@ -162,9 +162,9 @@ PRIVATE void benchmark_pagerank(graph_t* graph, void* rank,
}

/**
* Runs Dijkstra benchmark
* Runs SSSP benchmark
*/
PRIVATE void benchmark_dijkstra(graph_t* graph, void* distance,
PRIVATE void benchmark_sssp(graph_t* graph, void* distance,
totem_attr_t* attr) {
CALL_SAFE(sssp_hybrid(get_random_src(graph),
reinterpret_cast<weight_t*>(distance)));
Expand Down Expand Up @@ -211,7 +211,7 @@ PRIVATE void benchmark_run() {

graph_t* graph = NULL;
CALL_SAFE(graph_initialize(options->graph_file,
(options->benchmark == BENCHMARK_DIJKSTRA),
(options->benchmark == BENCHMARK_SSSP),
&graph));
print_config(graph, options, BENCHMARKS[options->benchmark].name);

Expand Down
16 changes: 8 additions & 8 deletions src/benchmark/totem_benchmark_cmdline.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,7 @@ PRIVATE void display_help(char* exe_name, int exit_err) {
" -bNUM Benchmark\n"
" %d: BFS (default)\n"
" %d: PageRank\n"
" %d: Dijkstra\n"
" %d: SSSP\n"
" %d: Betweenness\n"
" %d: Graph500\n"
" %d: Clustering Coefficient\n"
Expand Down Expand Up @@ -82,14 +82,14 @@ PRIVATE void display_help(char* exe_name, int exit_err) {
" %d: guided\n"
" -tNUM [1-%d] Number of CPU threads to use (default %d).\n"
" -h Print this help message\n",
exe_name, BENCHMARK_BFS, BENCHMARK_PAGERANK, BENCHMARK_DIJKSTRA,
BENCHMARK_BETWEENNESS, BENCHMARK_GRAPH500,
BENCHMARK_CLUSTERING_COEFFICIENT, get_gpu_count(), PAR_RANDOM,
PAR_SORTED_ASC, PAR_SORTED_DSC, GPU_GRAPH_MEM_DEVICE,
GPU_GRAPH_MEM_MAPPED, GPU_GRAPH_MEM_MAPPED_VERTICES,
exe_name, BENCHMARK_BFS, BENCHMARK_PAGERANK, BENCHMARK_SSSP,
BENCHMARK_BETWEENNESS, BENCHMARK_GRAPH500,
BENCHMARK_CLUSTERING_COEFFICIENT, get_gpu_count(), PAR_RANDOM,
PAR_SORTED_ASC, PAR_SORTED_DSC, GPU_GRAPH_MEM_DEVICE,
GPU_GRAPH_MEM_MAPPED, GPU_GRAPH_MEM_MAPPED_VERTICES,
GPU_GRAPH_MEM_MAPPED_EDGES, GPU_GRAPH_MEM_PARTITIONED_EDGES,
PLATFORM_CPU, PLATFORM_GPU, PLATFORM_HYBRID, REPEAT_MAX,
omp_sched_static, omp_sched_dynamic, omp_sched_guided,
PLATFORM_CPU, PLATFORM_GPU, PLATFORM_HYBRID, REPEAT_MAX,
omp_sched_static, omp_sched_dynamic, omp_sched_guided,
omp_get_max_threads(), omp_get_max_threads());
exit(exit_err);
}
Expand Down
Loading

0 comments on commit 4b10dc9

Please sign in to comment.