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/NN_kernel.cu | 145 ++++++++++++++++++++++++++++++++++++++++
 1 file changed, 145 insertions(+)
 create mode 100644 benchmarks/CUDA/NN/NN_kernel.cu

(limited to 'benchmarks/CUDA/NN/NN_kernel.cu')

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_
-- 
cgit v1.3