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/bfs.cu | 237 --------------------------------------------- 1 file changed, 237 deletions(-) delete mode 100644 benchmarks/CUDA/BFS/bfs.cu (limited to 'benchmarks/CUDA/BFS/bfs.cu') 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