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/NN/Makefile                    |  48 ++++
 benchmarks/CUDA/NN/NN.cu                       | 383 +++++++++++++++++++++++++
 benchmarks/CUDA/NN/NN_kernel.cu                | 145 ++++++++++
 benchmarks/CUDA/NN/README.GPGPU-Sim            |   2 +
 benchmarks/CUDA/NN/data/lw1.wei                | Bin 0 -> 624 bytes
 benchmarks/CUDA/NN/data/lw2.wei                | Bin 0 -> 31200 bytes
 benchmarks/CUDA/NN/data/lw3.wei                | Bin 0 -> 500400 bytes
 benchmarks/CUDA/NN/data/lw4.wei                | Bin 0 -> 4040 bytes
 benchmarks/CUDA/NN/data/t10k-images-idx3-ubyte | Bin 0 -> 7840016 bytes
 9 files changed, 578 insertions(+)
 create mode 100644 benchmarks/CUDA/NN/Makefile
 create mode 100644 benchmarks/CUDA/NN/NN.cu
 create mode 100644 benchmarks/CUDA/NN/NN_kernel.cu
 create mode 100644 benchmarks/CUDA/NN/README.GPGPU-Sim
 create mode 100644 benchmarks/CUDA/NN/data/lw1.wei
 create mode 100644 benchmarks/CUDA/NN/data/lw2.wei
 create mode 100644 benchmarks/CUDA/NN/data/lw3.wei
 create mode 100644 benchmarks/CUDA/NN/data/lw4.wei
 create mode 100644 benchmarks/CUDA/NN/data/t10k-images-idx3-ubyte

(limited to 'benchmarks/CUDA/NN')

diff --git a/benchmarks/CUDA/NN/Makefile b/benchmarks/CUDA/NN/Makefile
new file mode 100644
index 0000000..0d36972
--- /dev/null
+++ b/benchmarks/CUDA/NN/Makefile
@@ -0,0 +1,48 @@
+################################################################################
+#
+# 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	:= neuralnet
+# Cuda source files (compiled with cudacc)
+CUFILES		:= NN.cu 
+# C/C++ source files (compiled with gcc / c++)
+CCFILES		:= 
+GPGPUSIM_ROOT   := ../../..
+
+
+###############################################################################
+# Rules and targets
+
+include ../../../common/common.mk
diff --git a/benchmarks/CUDA/NN/NN.cu b/benchmarks/CUDA/NN/NN.cu
new file mode 100644
index 0000000..3ba5564
--- /dev/null
+++ b/benchmarks/CUDA/NN/NN.cu
@@ -0,0 +1,383 @@
+// includes, system
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+#include <math.h>
+#include <time.h>
+
+// includes, project
+#include <cutil.h>
+
+//#define NUM 10
+// includes, kernels
+#include <NN_kernel.cu>
+
+////////////////////////////////////////////////////////////////////////////////
+// declaration, forward
+
+extern "C"
+void computeGold(float*, const float*, const float*, unsigned int, unsigned int, unsigned int);
+void NeuralNetwork();
+
+unsigned g_verbose;
+unsigned NUM;
+////////////////////////////////////////////////////////////////////////////////
+// Program main
+////////////////////////////////////////////////////////////////////////////////
+int
+main(int argc, char** argv)
+{
+	int i, commandline_error;
+	commandline_error = 0;
+	g_verbose = 0;
+	if (argc >= 2) {
+		NUM = atoi(argv[1]);
+		for (i=2; i < argc;i++) {
+			if (argv[i][0] == '-') {
+				switch (argv[i][1]) {
+				case 'v': g_verbose = 1;
+					break;
+				default: commandline_error=1;
+				}
+			}
+			else commandline_error=1;
+		}
+	} else commandline_error=1;
+
+	if (commandline_error || !NUM) {
+		printf("Usage: ./NN <NUM> [-v]\n");
+		printf("where NUM is the number of images to process in parallel (up to 10000 for the t10k-images-idx3-ubyte database file) and -v is used to display approximately what each image looks like.\n");
+		return 1;
+	}
+
+	NeuralNetwork();
+    //CUT_EXIT(argc, argv);
+}
+
+void InitGPUMem(float *Layer1_Neurons_GPU,float *Layer1_Weights_GPU,float *Layer2_Neurons_GPU,float *Layer2_Weights_GPU,float *Layer3_Neurons_GPU,float *Layer3_Weights_GPU,float *Layer4_Neurons_GPU,float *Layer4_Weights_GPU,float *Layer5_Neurons_GPU)
+{
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer1_Neurons_GPU, sizeof(float)*29*29*NUM));
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer1_Weights_GPU, sizeof(float)*156));
+	
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer2_Neurons_GPU, sizeof(float)*13*13*6*NUM));
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer2_Weights_GPU, sizeof(float)*7800));
+
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer3_Neurons_GPU, sizeof(float)*1250*NUM));
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer3_Weights_GPU, sizeof(float)*125100));
+
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer4_Neurons_GPU, sizeof(float)*100*NUM));
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer4_Weights_GPU, sizeof(float)*1010));
+
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer5_Neurons_GPU, sizeof(float)*10*NUM));
+}
+void InitHostMem(float *Layer1_Weights_CPU,float *Layer2_Weights_CPU,float *Layer3_Weights_CPU,float *Layer4_Weights_CPU)
+{
+	// initial layer 1 weight
+	FILE * pFile1 = fopen ("data/lw1.wei","rb");
+	if (pFile1 != NULL)
+	{
+		for(int i=0;i<156;++i){
+			fread(&(Layer1_Weights_CPU[i]),sizeof(float),1,pFile1);
+			//printf("Layer1_Weights_CPU[%d]=%f\n", i, Layer1_Weights_CPU[i]);
+		}
+		fclose (pFile1);
+	}
+
+	// initial layer 2 weight
+	FILE * pFile2 = fopen ("data/lw2.wei","rb");
+	if (pFile2 != NULL)
+	{
+		fread(Layer2_Weights_CPU,sizeof(float),7800,pFile2);
+		fclose (pFile2);
+	}
+	// initial layer 3 weight
+	FILE * pFile3 = fopen ("data/lw3.wei","rb");
+	if (pFile3 != NULL)
+	{
+		fread(Layer3_Weights_CPU,sizeof(float),125100,pFile3);
+		fclose (pFile3);
+	}
+	// initial layer 4 weight
+	FILE * pFile4 = fopen ("data/lw4.wei","rb");
+	if (pFile4 != NULL)
+	{
+		fread(Layer4_Weights_CPU,sizeof(float),1010,pFile4);
+		fclose (pFile4);
+	}
+	if (!(pFile1 && pFile2 && pFile3 && pFile4))
+	{
+		printf("FAIL! INPUT WEIGHTS NOT FOUND!\n");
+		exit(1);
+	}
+}
+
+int swapEndianInt( int bEnum )
+{
+
+    int lEnum;
+    char *lE = (char*) &lEnum;
+    char *bE = (char*) &bEnum;
+    lE[0] = bE[3];
+    lE[1] = bE[2];
+    lE[2] = bE[1];
+    lE[3] = bE[0];
+    return lEnum;
+
+}
+void readIn(float *layer1)
+{
+	FILE *fp;
+	unsigned int *foo;
+	unsigned int i,j;
+	foo = (unsigned int *) calloc(sizeof(unsigned int),1);
+	//unsigned char image[29*29*NUM];
+	unsigned char* image = (unsigned char*) malloc(29*29*NUM * sizeof(char));
+	for (i=0;i<(29*29*NUM);i++) image[i]=0;
+	fp=fopen("data/t10k-images-idx3-ubyte","rt");
+	//fp=fopen("in.neu","rb");
+	if(fp)
+	{
+		fread(foo,sizeof(int),1,fp);
+		printf("magic number = %d\n", swapEndianInt(foo[0]));
+		fread(foo,sizeof(int),1,fp);
+		printf("number of items = %d\n", swapEndianInt(foo[0]));
+		fread(foo,sizeof(int),1,fp);
+		printf("number of rows = %d\n", swapEndianInt(foo[0]));
+		fread(foo,sizeof(int),1,fp);
+        printf("number of rows = %d\n", swapEndianInt(foo[0]));
+		for (j=0;j<NUM;j++) {
+			for (i=0;i<28;i++) 
+				fread((image+i*29+j*29*29),sizeof(char),28,fp);
+		}
+		//fread(layer1,sizeof(float),29*29,fp);
+		fclose(fp);
+		for (i=0;i<(29*29*NUM);i++)
+			layer1[i] = (1.0 - (float) image[i]/256);
+	}
+	else
+	{	
+		printf("FAIL! data/t10k-images-idx3-ubyte NOT FOUND!\n");
+		exit(1);
+	}
+}
+
+void output(double *final)
+{
+	int i;
+	FILE *fp=0;
+	fp=fopen("out.res","wb");
+	if(fp)
+	{
+		//for(i=0;i<10;i++) { 
+		//	printf("output[%d] = %e\n", i, final[i]);
+		//}
+		fwrite(final,sizeof(double),10,fp);
+		fclose(fp);
+	}
+}
+
+void NeuralNetwork()
+{	
+	int x,y;
+	// initialise card and timer
+	int deviceCount;                                                         
+	CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceCount(&deviceCount));                
+	if (deviceCount == 0) {                                                  
+		fprintf(stderr, "There is no device.\n");                            
+		exit(EXIT_FAILURE);                                                  
+	}                                                                        
+	int dev;                                                                 
+	for (dev = 0; dev < deviceCount; ++dev) {                                
+		cudaDeviceProp deviceProp;                                           
+		CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceProperties(&deviceProp, dev));   
+		if (deviceProp.major >= 1)                                           
+			break;                                                           
+	}                                                                        
+	if (dev == deviceCount) {                                                
+		fprintf(stderr, "There is no device supporting CUDA.\n");            
+		exit(EXIT_FAILURE);                                                  
+	}                                                                        
+	else                                                                     
+		CUDA_SAFE_CALL(cudaSetDevice(dev));  
+	//float Layer1_Neurons_CPU[29*29*NUM];
+	float *Layer1_Neurons_CPU = (float*) malloc (29*29*NUM * sizeof(float));
+	 /*={
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,0,0,0,0,0,0,0,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,0,0,0,1,1,0,0,0,0,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,0,0,1,1,1,1,1,1,0,0,0,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,0,0,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,0,0,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,0,0,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,0,0,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,0,0,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,0,0,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,0,0,0,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,0,0,0,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,0,0,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,0,0,0,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,0,0,0,0,0,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,0,0,0,0,0,0,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,0,0,0,0,0,0,0,0,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,0,0,0,0,0,0,0,0,0,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,0,0,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,
+1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1,1};*/
+
+	readIn(Layer1_Neurons_CPU);
+	if (g_verbose) {
+		for(y=0;y< 29*NUM;y++) {
+			if(!(y%29)) printf("\n");
+			for (x=0;x<29;x++) {
+				if (Layer1_Neurons_CPU[y*29+x]<0.5) {
+					printf("0");
+				}
+				else printf(" ");
+				//printf("%d", (Layer1_Neurons_CPU[y*29+x]>0.5));
+			}
+			printf("\n");
+			
+		}
+	}
+	float *Layer1_Neurons_GPU;
+	float Layer1_Weights_CPU[156];
+	float *Layer1_Weights_GPU;
+
+	float Layer2_Weights_CPU[7800];
+	float *Layer2_Weights_GPU;
+	float *Layer2_Neurons_GPU;
+
+	float Layer3_Weights_CPU[125100];
+	float *Layer3_Weights_GPU;
+	float *Layer3_Neurons_GPU;
+
+	float Layer4_Weights_CPU[1010];
+	float *Layer4_Weights_GPU;
+	float *Layer4_Neurons_GPU;
+
+	//float Layer5_Neurons_CPU[10*NUM];//={0,0,0,0,0,0,0,0,0,0};
+	float *Layer5_Neurons_CPU = (float*) malloc(10*NUM * sizeof(float));
+	for (x=0;x<10*NUM;x++) Layer5_Neurons_CPU[x]=0;
+	float *Layer5_Neurons_GPU;
+
+	double *outputLayer;
+	//unsigned int timer = 0;
+	//float totaltime = 0.0f;
+	//init input here
+	InitHostMem(Layer1_Weights_CPU,Layer2_Weights_CPU,Layer3_Weights_CPU,Layer4_Weights_CPU);
+
+
+	//allocate momory on Device
+	//InitGPUMem(Layer1_Neurons_GPU,Layer1_Weights_GPU,Layer2_Neurons_GPU,Layer2_Weights_GPU,Layer3_Neurons_GPU,Layer3_Weights_GPU,Layer4_Neurons_GPU,Layer4_Weights_GPU,Layer5_Neurons_GPU);
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer1_Neurons_GPU, sizeof(float)*29*29*NUM));
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer1_Weights_GPU, sizeof(float)*156));
+	
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer2_Neurons_GPU, sizeof(float)*13*13*6*NUM));
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer2_Weights_GPU, sizeof(float)*7800));
+
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer3_Neurons_GPU, sizeof(float)*1250*NUM));
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer3_Weights_GPU, sizeof(float)*125100));
+
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer4_Neurons_GPU, sizeof(float)*100*NUM));
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer4_Weights_GPU, sizeof(float)*1010));
+
+	CUDA_SAFE_CALL(cudaMalloc((void**) &Layer5_Neurons_GPU, sizeof(float)*10*NUM));
+	outputLayer = (double*)malloc(sizeof(double)*10*NUM);
+	//init 29x29 handwritting array
+	// already done in "initial"
+
+	//copy from CPU to GPU
+	CUDA_SAFE_CALL(cudaMemcpy(Layer1_Neurons_GPU,Layer1_Neurons_CPU, sizeof(float)*29*29*NUM, cudaMemcpyHostToDevice));
+	CUDA_SAFE_CALL(cudaMemcpy(Layer1_Weights_GPU,Layer1_Weights_CPU, sizeof(float)*156, cudaMemcpyHostToDevice));
+	CUDA_SAFE_CALL(cudaMemcpy(Layer2_Weights_GPU,Layer2_Weights_CPU, sizeof(float)*7800, cudaMemcpyHostToDevice));
+	CUDA_SAFE_CALL(cudaMemcpy(Layer3_Weights_GPU,Layer3_Weights_CPU, sizeof(float)*125100, cudaMemcpyHostToDevice));
+	CUDA_SAFE_CALL(cudaMemcpy(Layer4_Weights_GPU,Layer4_Weights_CPU, sizeof(float)*1010, cudaMemcpyHostToDevice));
+	CUDA_SAFE_CALL(cudaMemcpy(Layer5_Neurons_GPU,Layer5_Neurons_CPU, sizeof(float)*10*NUM, cudaMemcpyHostToDevice));
+
+	// CUT_SAFE_CALL(cutCreateTimer(&timer));
+	// CUT_SAFE_CALL(cutStartTimer(timer));
+	printf("NUM=%d\n", NUM);
+	dim3 Layer1_Block(6,NUM,1);
+	dim3 Layer1_Thread(13,13);
+	executeFirstLayer<<<Layer1_Block,Layer1_Thread>>>(Layer1_Neurons_GPU,Layer1_Weights_GPU,Layer2_Neurons_GPU);
+
+	dim3 Layer2_Block(50,NUM,1);
+	dim3 Layer2_Thread(5,5);
+	executeSecondLayer<<<Layer2_Block,Layer2_Thread>>>(Layer2_Neurons_GPU, Layer2_Weights_GPU,Layer3_Neurons_GPU);
+
+	dim3 Layer3_Block(100,NUM,1);
+	dim3 Layer3_Thread(1,1);
+	executeThirdLayer<<<Layer3_Block,Layer3_Thread>>>(Layer3_Neurons_GPU, Layer3_Weights_GPU,Layer4_Neurons_GPU);
+
+	dim3 Layer4_Block(10,NUM,1);
+	dim3 Layer4_Thread(1,1);
+	executeFourthLayer<<<Layer4_Block,Layer4_Thread>>>(Layer4_Neurons_GPU,Layer4_Weights_GPU,Layer5_Neurons_GPU);
+
+	CUT_CHECK_ERROR("Kernel execution failed");
+
+  //  CUT_SAFE_CALL(cutStopTimer(timer));
+
+//	totaltime = cutGetTimerValue(timer);
+
+	//copy from GPU to CPU
+    CUDA_SAFE_CALL(cudaMemcpy(Layer5_Neurons_CPU,Layer5_Neurons_GPU, sizeof(float)*10*NUM, cudaMemcpyDeviceToHost));
+    
+    // stop and destroy timer
+
+    //printf("Processing time: %f (ms) \n", totaltime);
+	//  CUT_SAFE_CALL(cutDeleteTimer(timer));
+
+	for(int a=0;a<10*NUM;a++)
+	{
+		//printf("output[%d]=%f\n", a, Layer5_Neurons_CPU[a]);
+		outputLayer[a] = (double)Layer5_Neurons_CPU[a];
+		if (!(a%10)) {
+			if (a) printf("%d ", y);
+			x=outputLayer[a];
+			y=0;
+		}
+		if (outputLayer[a]>x) {
+			x=outputLayer[a];
+			y=a%10;
+		}	
+	}
+	printf("%d\n", y);
+	output(outputLayer);
+
+	/*
+	//float Layer4_Neurons_CPU[100*NUM];
+	float *Layer4_Neurons_CPU = (float*) malloc(100*NUM*sizeof(float));
+	CUDA_SAFE_CALL(cudaMemcpy(Layer4_Neurons_CPU,Layer4_Neurons_GPU,sizeof(float)*100,cudaMemcpyDeviceToHost));
+	FILE *fp=fopen("layer_4.neu","wb");
+	fwrite(Layer4_Neurons_CPU,sizeof(float),100*NUM,fp);
+	fclose(fp);
+
+	//float Layer3_Neurons_CPU[50*5*5*NUM];
+	float *Layer3_Neurons_CPU = (float*) malloc(50*5*5*NUM*sizeof(float));
+	CUDA_SAFE_CALL(cudaMemcpy(Layer3_Neurons_CPU,Layer3_Neurons_GPU,sizeof(float)*50*5*5,cudaMemcpyDeviceToHost));
+	fp=fopen("layer_3.neu","wb");
+	fwrite(Layer3_Neurons_CPU,sizeof(float),50*5*5*NUM,fp);
+	fclose(fp);
+
+	//float Layer2_Neurons_CPU[13*13*6*NUM];
+	float *Layer2_Neurons_CPU = (float*) malloc(13*13*6*NUM*sizeof(float));
+	CUDA_SAFE_CALL(cudaMemcpy(Layer2_Neurons_CPU,Layer2_Neurons_GPU,sizeof(float)*13*13*6,cudaMemcpyDeviceToHost));
+	fp=fopen("layer_2.neu","wb");
+	fwrite(Layer2_Neurons_CPU,sizeof(float),13*13*6*NUM,fp);
+	fclose(fp);
+
+	fp=fopen("layer_1.neu","wb");
+	fwrite(Layer1_Neurons_CPU,sizeof(float),29*29*NUM,fp);
+	fclose(fp);    */
+
+	exit(0);
+}
+
diff --git a/benchmarks/CUDA/NN/NN_kernel.cu b/benchmarks/CUDA/NN/NN_kernel.cu
new file mode 100644
index 0000000..030319d
--- /dev/null
+++ b/benchmarks/CUDA/NN/NN_kernel.cu
@@ -0,0 +1,145 @@
+#ifndef _NN_KERNEL_H_
+#define _NN_KERNEL_H_
+
+#include <stdio.h>
+
+#define CHECK_BANK_CONFLICTS 0
+#if CHECK_BANK_CONFLICTS
+#define AS(i, j) CUT_BANK_CHECKER(((float*)&As[0][0]), (BLOCK_SIZE * i + j))
+#define BS(i, j) CUT_BANK_CHECKER(((float*)&Bs[0][0]), (BLOCK_SIZE * i + j))
+#else
+#define AS(i, j) As[i][j]
+#define BS(i, j) Bs[i][j]
+#endif
+
+
+__constant__ int kernelTemplate[25] = {
+        0,  1,  2,  3,  4,
+        29, 30, 31, 32, 33,
+        58, 59, 60, 61, 62,
+        87, 88, 89, 90, 91,
+        116,117,118,119,120 };
+
+__global__ void executeFirstLayer(float *Layer1_Neurons_GPU,float *Layer1_Weights_GPU,float *Layer2_Neurons_GPU)
+{
+	int blockID=blockIdx.x;
+	int pixelX=threadIdx.x;
+	int pixelY=threadIdx.y;
+
+
+	int weightBegin=blockID*26;
+	int windowX=pixelX*2;
+	int windowY=pixelY*2;
+
+	float result=0;
+
+	result+=Layer1_Weights_GPU[weightBegin];
+
+	++weightBegin;
+
+	for(int i=0;i<25;++i)
+	{
+		result+=Layer1_Neurons_GPU[(windowY*29+windowX+kernelTemplate[i])+(29*29*blockIdx.y)]*Layer1_Weights_GPU[weightBegin+i];
+	}
+
+	result=(1.7159*tanhf(0.66666667*result));
+
+	Layer2_Neurons_GPU[(13*13*blockID+pixelY*13+pixelX)+(13*13*6*blockIdx.y)]=result;
+
+}
+
+__constant__ int kernelTemplate2[25] = {
+        0,  1,  2,  3,  4,
+        13, 14, 15, 16, 17, 
+        26, 27, 28, 29, 30,
+        39, 40, 41, 42, 43, 
+        52, 53, 54, 55, 56   };
+
+__global__ void executeSecondLayer(float *Layer2_Neurons_GPU, float *Layer2_Weights_GPU,float *Layer3_Neurons_GPU)
+{
+	int blockID=blockIdx.x;
+	int pixelX=threadIdx.x;
+	int pixelY=threadIdx.y;
+
+
+	int weightBegin=blockID*26*6;
+	int windowX=pixelX*2;
+	int windowY=pixelY*2;
+    
+	float result=0;
+
+	
+	result+=Layer2_Weights_GPU[weightBegin];
+	
+	if(blockID==1 && pixelX==0 && pixelY==0)
+	{
+		result+=0;
+	}
+
+	++weightBegin;
+
+	for (int i=0; i<25; ++i )
+    {
+        result+=Layer2_Neurons_GPU[(windowX + 13*windowY +kernelTemplate2[i])+(13*13*6*blockIdx.y)]*Layer2_Weights_GPU[weightBegin+i*6];
+        result+=Layer2_Neurons_GPU[(169 + windowX + 13*windowY +kernelTemplate2[i])+(13*13*6*blockIdx.y)]*Layer2_Weights_GPU[weightBegin+i*6+1];
+	result+=Layer2_Neurons_GPU[(338 + windowX + 13*windowY + kernelTemplate2[i])+(13*13*6*blockIdx.y)]*Layer2_Weights_GPU[weightBegin+i*6+2];
+        result+=Layer2_Neurons_GPU[(507 + windowX + 13*windowY + kernelTemplate2[i])+(13*13*6*blockIdx.y)]*Layer2_Weights_GPU[weightBegin+i*6+3];
+        result+=Layer2_Neurons_GPU[(676 + windowX + 13*windowY + kernelTemplate2[i])+(13*13*6*blockIdx.y)]*Layer2_Weights_GPU[weightBegin+i*6+4];
+        result+=Layer2_Neurons_GPU[(845 + windowX + 13*windowY + kernelTemplate2[i])+(13*13*6*blockIdx.y)]*Layer2_Weights_GPU[weightBegin+i*6+5];
+	}
+
+	result=(1.7159*tanhf(0.66666667*result));
+
+	Layer3_Neurons_GPU[(5*5*blockID+pixelY*5+pixelX)+(1250*blockIdx.y)]=result;
+}
+
+__global__ void executeThirdLayer(float *Layer3_Neurons_GPU, float *Layer3_Weights_GPU,float *Layer4_Neurons_GPU)
+{
+	int blockID=blockIdx.x;
+	//int pixelY=threadIdx.y;
+
+
+	int weightBegin=blockID*1251;
+ 
+	float result=0;
+
+	result+=Layer3_Weights_GPU[weightBegin];
+
+	++weightBegin;
+
+    for (int i=0; i<1250; ++i )
+    {
+		result+=Layer3_Neurons_GPU[i+(1250*blockIdx.y)]*Layer3_Weights_GPU[weightBegin+i];
+    }
+
+	result=(1.7159*tanhf(0.66666667*result));
+
+	Layer4_Neurons_GPU[blockID+(100*blockIdx.y)]=result;
+
+}
+
+__global__ void executeFourthLayer(float *Layer4_Neurons_GPU,float *Layer4_Weights_GPU,float *Layer5_Neurons_GPU)
+{
+	int blockID=blockIdx.x;
+	//int pixelY=threadIdx.y;
+
+
+	int weightBegin=blockID*101;
+ 
+	float result=0;
+
+	result+=Layer4_Weights_GPU[weightBegin];
+
+	++weightBegin;
+
+    for (int i=0; i<100; ++i )
+    {
+		result+=Layer4_Neurons_GPU[i+(100*blockIdx.y)]*Layer4_Weights_GPU[weightBegin+i];
+    }
+
+	result=(1.7159*tanhf(0.66666667*result));
+
+	Layer5_Neurons_GPU[blockID+(10*blockIdx.y)]=result;
+}
+
+#endif // #ifndef _NN_KERNEL_H_
diff --git a/benchmarks/CUDA/NN/README.GPGPU-Sim b/benchmarks/CUDA/NN/README.GPGPU-Sim
new file mode 100644
index 0000000..a55ac20
--- /dev/null
+++ b/benchmarks/CUDA/NN/README.GPGPU-Sim
@@ -0,0 +1,2 @@
+make
+./gpgpu_ptx_sim__neuralnet 28
diff --git a/benchmarks/CUDA/NN/data/lw1.wei b/benchmarks/CUDA/NN/data/lw1.wei
new file mode 100644
index 0000000..d6e140e
Binary files /dev/null and b/benchmarks/CUDA/NN/data/lw1.wei differ
diff --git a/benchmarks/CUDA/NN/data/lw2.wei b/benchmarks/CUDA/NN/data/lw2.wei
new file mode 100644
index 0000000..1a340ea
Binary files /dev/null and b/benchmarks/CUDA/NN/data/lw2.wei differ
diff --git a/benchmarks/CUDA/NN/data/lw3.wei b/benchmarks/CUDA/NN/data/lw3.wei
new file mode 100644
index 0000000..e86a6b6
Binary files /dev/null and b/benchmarks/CUDA/NN/data/lw3.wei differ
diff --git a/benchmarks/CUDA/NN/data/lw4.wei b/benchmarks/CUDA/NN/data/lw4.wei
new file mode 100644
index 0000000..17cf469
Binary files /dev/null and b/benchmarks/CUDA/NN/data/lw4.wei differ
diff --git a/benchmarks/CUDA/NN/data/t10k-images-idx3-ubyte b/benchmarks/CUDA/NN/data/t10k-images-idx3-ubyte
new file mode 100644
index 0000000..1170b2c
Binary files /dev/null and b/benchmarks/CUDA/NN/data/t10k-images-idx3-ubyte differ
-- 
cgit v1.3