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]

1 files changed, 0 insertions, 237 deletions

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 <stdlib.h>

-#include <stdio.h>

-#include <string.h>

-#include <math.h>

-#include <cutil.h>

-

-#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 <kernel.cu>

-

-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<no_of_nodes;i++)

-	h_cost[i]=-1;

-	h_cost[source]=0;

-

-	// allocate device memory for result

-    int* d_cost;

-    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_cost, sizeof(int)*no_of_nodes));

-    CUDA_SAFE_CALL( cudaMemcpy( d_cost, h_cost, sizeof(int)*no_of_nodes, cudaMemcpyHostToDevice) );

-

-    //make a bool to check if the execution is over

-    bool *d_over;

-    CUDA_SAFE_CALL( cudaMalloc( (void**) &d_over, sizeof(bool)));

-

-	printf("Copied Everything to GPU memory\n");

-    

-    // setup execution parameters

-    dim3  grid( num_of_blocks, 1, 1);

-    dim3  threads( num_of_threads_per_block, 1, 1);

-

-	//start the timer

-    unsigned int timer = 0;

-    float timer_acc = 0.0f;

-    CUT_SAFE_CALL( cutCreateTimer( &timer));

-    CUT_SAFE_CALL( cutStartTimer( timer));

-	int k=0;

-	

-	bool stop;

-	//Call the Kernel untill all the elements of Frontier are not false

-    do

-    {

-	//if no thread changes this value then the loop stops

-	stop=false;

-	CUDA_SAFE_CALL( cudaMemcpy( d_over, &stop, sizeof(bool), cudaMemcpyHostToDevice) );

-    	CUT_SAFE_CALL( cutStartTimer( timer));

-	Kernel<<< grid, threads, 0 >>>( 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<no_of_nodes;i++)

-	fprintf(fpo,"%d) cost:%d\n",i,h_cost[i]);

-	fclose(fpo);

-	printf("Result stored in result.txt\n");

-	

-	

-    // cleanup memory

-    free( h_graph_nodes);

-    free( h_graph_edges);

-    free( h_graph_mask);

-    free( h_graph_visited);

-    free( h_cost);

-    CUDA_SAFE_CALL(cudaFree(d_graph_nodes));

-    CUDA_SAFE_CALL(cudaFree(d_graph_edges));

-    CUDA_SAFE_CALL(cudaFree(d_graph_mask));

-    CUDA_SAFE_CALL(cudaFree(d_graph_visited));

-    CUDA_SAFE_CALL(cudaFree(d_cost));

-}