From 69f2911e04ffb1b19eef1fafb8c040af271f656e Mon Sep 17 00:00:00 2001
From: Tor Aamodt <aamodt@ece.ubc.ca>
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/kernel.cu | 43 +++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 43 insertions(+)
 create mode 100644 benchmarks/CUDA/BFS/kernel.cu

(limited to 'benchmarks/CUDA/BFS/kernel.cu')

diff --git a/benchmarks/CUDA/BFS/kernel.cu b/benchmarks/CUDA/BFS/kernel.cu
new file mode 100644
index 0000000..50f9e20
--- /dev/null
+++ b/benchmarks/CUDA/BFS/kernel.cu
@@ -0,0 +1,43 @@
+/*********************************************************************************
+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.
+
+The CUDA Kernel for Applying BFS on a loaded Graph. Created By Pawan Harish
+**********************************************************************************/
+#ifndef _KERNEL_H_
+#define _KERNEL_H_
+
+__global__ void
+Kernel( Node* g_graph_nodes, int* g_graph_edges, bool* g_graph_mask, bool* g_graph_visited, int* g_cost, bool *g_over, int no_of_nodes) 
+{
+	int tid = blockIdx.x*MAX_THREADS_PER_BLOCK + threadIdx.x;
+	if( tid<no_of_nodes && g_graph_mask[tid])
+	{
+		g_graph_mask[tid]=false;
+		g_graph_visited[tid]=true;
+		for(int i=g_graph_nodes[tid].starting; i<(g_graph_nodes[tid].no_of_edges + g_graph_nodes[tid].starting); i++)
+			{
+			int id = g_graph_edges[i];
+			if(!g_graph_visited[id])
+				{
+				g_cost[id]=g_cost[tid]+1;
+				g_graph_mask[id]=true;
+				//Change the loop stop value such that loop continues
+				*g_over=true;
+				}
+			}
+	}
+}
+
+#endif 
-- 
cgit v1.3