From 69f2911e04ffb1b19eef1fafb8c040af271f656e Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Thu, 15 Jul 2010 18:09:46 -0800 Subject: creating branch for adding support for CUDA 3.x and Fermi [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829] --- benchmarks/CUDA/BFS/Makefile | 49 + benchmarks/CUDA/BFS/README.GPGPU-Sim | 4 + benchmarks/CUDA/BFS/ReadMe.txt | 25 + benchmarks/CUDA/BFS/SampleGraph.jpg | Bin 0 -> 59715 bytes benchmarks/CUDA/BFS/bfs.cu | 237 + benchmarks/CUDA/BFS/data/GraphFormat.txt | 20 + benchmarks/CUDA/BFS/data/SampleGraph.txt | 31 + benchmarks/CUDA/BFS/data/graph4096.txt | 28677 ++ benchmarks/CUDA/BFS/data/graph65536.txt | 458757 +++++++++++++++++++++ benchmarks/CUDA/BFS/data/result_SampleGraph.txt | 6 + benchmarks/CUDA/BFS/kernel.cu | 43 + 11 files changed, 487849 insertions(+) create mode 100644 benchmarks/CUDA/BFS/Makefile create mode 100644 benchmarks/CUDA/BFS/README.GPGPU-Sim create mode 100644 benchmarks/CUDA/BFS/ReadMe.txt create mode 100644 benchmarks/CUDA/BFS/SampleGraph.jpg create mode 100644 benchmarks/CUDA/BFS/bfs.cu create mode 100644 benchmarks/CUDA/BFS/data/GraphFormat.txt create mode 100644 benchmarks/CUDA/BFS/data/SampleGraph.txt create mode 100644 benchmarks/CUDA/BFS/data/graph4096.txt create mode 100644 benchmarks/CUDA/BFS/data/graph65536.txt create mode 100644 benchmarks/CUDA/BFS/data/result_SampleGraph.txt create mode 100644 benchmarks/CUDA/BFS/kernel.cu (limited to 'benchmarks/CUDA/BFS') diff --git a/benchmarks/CUDA/BFS/Makefile b/benchmarks/CUDA/BFS/Makefile new file mode 100644 index 0000000..76d8b17 --- /dev/null +++ b/benchmarks/CUDA/BFS/Makefile @@ -0,0 +1,49 @@ +################################################################################ +# +# Copyright 1993-2006 NVIDIA Corporation. All rights reserved. +# +# NOTICE TO USER: +# +# This source code is subject to NVIDIA ownership rights under U.S. and +# international Copyright laws. +# +# NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE +# CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR +# IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH +# REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF +# MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. +# IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, +# OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS +# OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE +# OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE +# OR PERFORMANCE OF THIS SOURCE CODE. +# +# U.S. Government End Users. This source code is a "commercial item" as +# that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of +# "commercial computer software" and "commercial computer software +# documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) +# and is provided to the U.S. Government only as a commercial end item. +# Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through +# 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the +# source code with only those rights set forth herein. +# +################################################################################ +# +# Build script for project +# +################################################################################ + +# Add source files here +EXECUTABLE := bfs +# Cuda source files (compiled with cudacc) +CUFILES := bfs.cu +# C/C++ source files (compiled with gcc / c++) +CCFILES := + +GPGPUSIM_ROOT := ../../.. + +################################################################################ +# Rules and targets + +include ../../../common/common.mk + diff --git a/benchmarks/CUDA/BFS/README.GPGPU-Sim b/benchmarks/CUDA/BFS/README.GPGPU-Sim new file mode 100644 index 0000000..a83f4cb --- /dev/null +++ b/benchmarks/CUDA/BFS/README.GPGPU-Sim @@ -0,0 +1,4 @@ +To build and run: + +make +./gpgpu_ptx_sim__bfs ./data/graph65536.txt diff --git a/benchmarks/CUDA/BFS/ReadMe.txt b/benchmarks/CUDA/BFS/ReadMe.txt new file mode 100644 index 0000000..71a9c03 --- /dev/null +++ b/benchmarks/CUDA/BFS/ReadMe.txt @@ -0,0 +1,25 @@ +CUDA Implementations of Breadth First Search and Single Source Shortest Path +algorithms given in IEEE High performance 2007 paper + + +"Accelerating Large Graph Algorithms on the GPU using CUDA" +By: Pawan Harish and P.J.Narayanan + + +Please note the given implementation of SSSP requires CUDA 1.1 Hardware whereas BFS does not +require this hardware. These codes are not optimized, please use CUDA profiler for +better optimization. + + +Copyright (c) 2008 International Institute of Information Technology. +All rights reserved. + +Permission to use, copy, modify and distribute this software and its documentation for +educational purpose is hereby granted without fee, provided that the above copyright +notice and this permission notice appear in all copies of this software and that you do +not sell the software. + +THE SOFTWARE IS PROVIDED "AS IS" AND WITHOUT WARRANTY OF ANY KIND,EXPRESS, IMPLIED OR +OTHERWISE. + +Created by Pawan Harish. \ No newline at end of file diff --git a/benchmarks/CUDA/BFS/SampleGraph.jpg b/benchmarks/CUDA/BFS/SampleGraph.jpg new file mode 100644 index 0000000..3d56305 Binary files /dev/null and b/benchmarks/CUDA/BFS/SampleGraph.jpg differ diff --git a/benchmarks/CUDA/BFS/bfs.cu b/benchmarks/CUDA/BFS/bfs.cu new file mode 100644 index 0000000..8e2348b --- /dev/null +++ b/benchmarks/CUDA/BFS/bfs.cu @@ -0,0 +1,237 @@ +/*********************************************************************************** +Implementing Breadth first search on CUDA using algorithm given in HiPC'07 + paper "Accelerating Large Graph Algorithms on the GPU using CUDA" + +Copyright (c) 2008 International Institute of Information Technology. +All rights reserved. + +Permission to use, copy, modify and distribute this software and its documentation for +educational purpose is hereby granted without fee, provided that the above copyright +notice and this permission notice appear in all copies of this software and that you do +not sell the software. + +THE SOFTWARE IS PROVIDED "AS IS" AND WITHOUT WARRANTY OF ANY KIND,EXPRESS, IMPLIED OR +OTHERWISE. + +Created by Pawan Harish. +************************************************************************************/ +#include +#include +#include +#include +#include + +#define MAX_THREADS_PER_BLOCK 256 + +int no_of_nodes; +int edge_list_size; +FILE *fp; + +//Structure to hold a node information +struct Node +{ +int starting; +int no_of_edges; +}; + +#include + +void BFSGraph(int argc, char** argv); + +//////////////////////////////////////////////////////////////////////////////// +// Main Program +//////////////////////////////////////////////////////////////////////////////// +int main( int argc, char** argv) +{ + no_of_nodes=0; + edge_list_size=0; + BFSGraph( argc, argv); + //CUT_EXIT(argc, argv); + return 0; +} + + + +//////////////////////////////////////////////////////////////////////////////// +//Apply BFS on a Graph using CUDA +//////////////////////////////////////////////////////////////////////////////// +void BFSGraph( int argc, char** argv) +{ + +// CUT_DEVICE_INIT(); + + printf("Reading File\n"); + static char *input_file_name; + //printf("argc=%d\n", argc); + if (argc == 2 ) { + input_file_name = argv[1]; + printf("Input file: %s\n", input_file_name); + } + else + { + input_file_name = "SampleGraph.txt"; + printf("No input file specified, defaulting to SampleGraph.txt\n"); + } + //Read in Graph from a file + fp = fopen(input_file_name,"r"); + if(!fp) + { + printf("Error Reading graph file\n"); + return; + } + + int source = 0; + + fscanf(fp,"%d",&no_of_nodes); + + int num_of_blocks = 1; + int num_of_threads_per_block = no_of_nodes; + + //Make execution Parameters according to the number of nodes + //Distribute threads across multiple Blocks if necessary + if(no_of_nodes>MAX_THREADS_PER_BLOCK) + { + num_of_blocks = (int)ceil(no_of_nodes/(double)MAX_THREADS_PER_BLOCK); + num_of_threads_per_block = MAX_THREADS_PER_BLOCK; + } + + // allocate host memory + Node* h_graph_nodes = (Node*) malloc(sizeof(Node)*no_of_nodes); + bool *h_graph_mask = (bool*) malloc(sizeof(bool)*no_of_nodes); + bool *h_graph_visited = (bool*) malloc(sizeof(bool)*no_of_nodes); + + int start, edgeno; + // initalize the memory + for( unsigned int i = 0; i < no_of_nodes; i++) + { + fscanf(fp,"%d %d",&start,&edgeno); + h_graph_nodes[i].starting = start; + h_graph_nodes[i].no_of_edges = edgeno; + h_graph_mask[i]=false; + h_graph_visited[i]=false; + } + + //read the source node from the file + fscanf(fp,"%d",&source); + + //set the source node as true in the mask + h_graph_mask[source]=true; + + fscanf(fp,"%d",&edge_list_size); + + int id,cost; + int* h_graph_edges = (int*) malloc(sizeof(int)*edge_list_size); + for(int i=0; i < edge_list_size ; i++) + { + fscanf(fp,"%d",&id); + fscanf(fp,"%d",&cost); + h_graph_edges[i] = id; + } + + if(fp) + fclose(fp); + + printf("Read File\n"); + + //Copy the Node list to device memory + Node* d_graph_nodes; + CUDA_SAFE_CALL( cudaMalloc( (void**) &d_graph_nodes, sizeof(Node)*no_of_nodes) ); + CUDA_SAFE_CALL( cudaMemcpy( d_graph_nodes, h_graph_nodes, sizeof(Node)*no_of_nodes, cudaMemcpyHostToDevice) ); + + //Copy the Edge List to device Memory + int* d_graph_edges; + CUDA_SAFE_CALL( cudaMalloc( (void**) &d_graph_edges, sizeof(int)*edge_list_size) ); + CUDA_SAFE_CALL( cudaMemcpy( d_graph_edges, h_graph_edges, sizeof(int)*edge_list_size, cudaMemcpyHostToDevice) ); + + //Copy the Mask to device memory + bool* d_graph_mask; + CUDA_SAFE_CALL( cudaMalloc( (void**) &d_graph_mask, sizeof(bool)*no_of_nodes) ); + CUDA_SAFE_CALL( cudaMemcpy( d_graph_mask, h_graph_mask, sizeof(bool)*no_of_nodes, cudaMemcpyHostToDevice) ); + + //Copy the Visited nodes array to device memory + bool* d_graph_visited; + CUDA_SAFE_CALL( cudaMalloc( (void**) &d_graph_visited, sizeof(bool)*no_of_nodes) ); + CUDA_SAFE_CALL( cudaMemcpy( d_graph_visited, h_graph_visited, sizeof(bool)*no_of_nodes, cudaMemcpyHostToDevice) ); + + // allocate mem for the result on host side + int* h_cost = (int*) malloc( sizeof(int)*no_of_nodes); + for(int i=0;i>>( d_graph_nodes, d_graph_edges, d_graph_mask, d_graph_visited, d_cost, d_over, no_of_nodes); + CUDA_SAFE_CALL(cudaThreadSynchronize()); + CUT_SAFE_CALL( cutStopTimer( timer)); + timer_acc += cutGetTimerValue(timer); + CUT_SAFE_CALL( cutResetTimer( timer)); + // check if kernel execution generated and error + CUT_CHECK_ERROR("Kernel execution failed"); + CUDA_SAFE_CALL( cudaMemcpy( &stop, d_over, sizeof(bool), cudaMemcpyDeviceToHost) ); + k++; + } + while(stop); + + + printf("Kernel Executed %d times\n",k); + + // copy result from device to host + CUDA_SAFE_CALL( cudaMemcpy( h_cost, d_cost, sizeof(int)*no_of_nodes, cudaMemcpyDeviceToHost) ); + + //Stop the Timer + CUT_SAFE_CALL( cutStopTimer( timer)); + //printf( "Processing time: %f (ms)\n", cutGetTimerValue(timer)); + + printf( "Processing time: %f (ms)\n", timer_acc); + CUT_SAFE_CALL( cutDeleteTimer( timer)); + + + //Store the result into a file + FILE *fpo = fopen("result.txt","w"); + for(int i=0;i