From 11b308e7363e937966b035b4891db32b4eece3bf Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Fri, 1 Oct 2010 08:55:28 -0800 Subject: integrating recent changes from fermi-test into fermi (i'll use "fermi" for more disruptive changes to the pipeline model such as updating the MSHRs and getting rid of the warp tracker, ripping out DWF, etc...) [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7805] --- benchmarks/CUDA/BFS/Makefile | 49 - benchmarks/CUDA/BFS/README.GPGPU-Sim | 4 - benchmarks/CUDA/BFS/ReadMe.txt | 25 - benchmarks/CUDA/BFS/SampleGraph.jpg | Bin 59715 -> 0 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 deletions(-) delete mode 100644 benchmarks/CUDA/BFS/Makefile delete mode 100644 benchmarks/CUDA/BFS/README.GPGPU-Sim delete mode 100644 benchmarks/CUDA/BFS/ReadMe.txt delete mode 100644 benchmarks/CUDA/BFS/SampleGraph.jpg delete mode 100644 benchmarks/CUDA/BFS/bfs.cu delete mode 100644 benchmarks/CUDA/BFS/data/GraphFormat.txt delete mode 100644 benchmarks/CUDA/BFS/data/SampleGraph.txt delete mode 100644 benchmarks/CUDA/BFS/data/graph4096.txt delete mode 100644 benchmarks/CUDA/BFS/data/graph65536.txt delete mode 100644 benchmarks/CUDA/BFS/data/result_SampleGraph.txt delete mode 100644 benchmarks/CUDA/BFS/kernel.cu (limited to 'benchmarks/CUDA/BFS') diff --git a/benchmarks/CUDA/BFS/Makefile b/benchmarks/CUDA/BFS/Makefile deleted file mode 100644 index 76d8b17..0000000 --- a/benchmarks/CUDA/BFS/Makefile +++ /dev/null @@ -1,49 +0,0 @@ -################################################################################ -# -# 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 deleted file mode 100644 index a83f4cb..0000000 --- a/benchmarks/CUDA/BFS/README.GPGPU-Sim +++ /dev/null @@ -1,4 +0,0 @@ -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 deleted file mode 100644 index 71a9c03..0000000 --- a/benchmarks/CUDA/BFS/ReadMe.txt +++ /dev/null @@ -1,25 +0,0 @@ -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 deleted file mode 100644 index 3d56305..0000000 Binary files a/benchmarks/CUDA/BFS/SampleGraph.jpg and /dev/null differ diff --git a/benchmarks/CUDA/BFS/bfs.cu b/benchmarks/CUDA/BFS/bfs.cu deleted file mode 100644 index 8e2348b..0000000 --- a/benchmarks/CUDA/BFS/bfs.cu +++ /dev/null @@ -1,237 +0,0 @@ -/*********************************************************************************** -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