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/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 624 -> 0 bytes benchmarks/CUDA/NN/data/lw2.wei | Bin 31200 -> 0 bytes benchmarks/CUDA/NN/data/lw3.wei | Bin 500400 -> 0 bytes benchmarks/CUDA/NN/data/lw4.wei | Bin 4040 -> 0 bytes benchmarks/CUDA/NN/data/t10k-images-idx3-ubyte | Bin 7840016 -> 0 bytes 9 files changed, 578 deletions(-) delete mode 100644 benchmarks/CUDA/NN/Makefile delete mode 100644 benchmarks/CUDA/NN/NN.cu delete mode 100644 benchmarks/CUDA/NN/NN_kernel.cu delete mode 100644 benchmarks/CUDA/NN/README.GPGPU-Sim delete mode 100644 benchmarks/CUDA/NN/data/lw1.wei delete mode 100644 benchmarks/CUDA/NN/data/lw2.wei delete mode 100644 benchmarks/CUDA/NN/data/lw3.wei delete mode 100644 benchmarks/CUDA/NN/data/lw4.wei delete 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 deleted file mode 100644 index 0d36972..0000000 --- a/benchmarks/CUDA/NN/Makefile +++ /dev/null @@ -1,48 +0,0 @@ -################################################################################ -# -# 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 deleted file mode 100644 index 3ba5564..0000000 --- a/benchmarks/CUDA/NN/NN.cu +++ /dev/null @@ -1,383 +0,0 @@ -// includes, system -#include -#include -#include -#include -#include - -// includes, project -#include - -//#define NUM 10 -// includes, kernels -#include - -//////////////////////////////////////////////////////////////////////////////// -// 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 [-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= 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_Neurons_GPU,Layer1_Weights_GPU,Layer2_Neurons_GPU); - - dim3 Layer2_Block(50,NUM,1); - dim3 Layer2_Thread(5,5); - executeSecondLayer<<>>(Layer2_Neurons_GPU, Layer2_Weights_GPU,Layer3_Neurons_GPU); - - dim3 Layer3_Block(100,NUM,1); - dim3 Layer3_Thread(1,1); - executeThirdLayer<<>>(Layer3_Neurons_GPU, Layer3_Weights_GPU,Layer4_Neurons_GPU); - - dim3 Layer4_Block(10,NUM,1); - dim3 Layer4_Thread(1,1); - executeFourthLayer<<>>(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 deleted file mode 100644 index 030319d..0000000 --- a/benchmarks/CUDA/NN/NN_kernel.cu +++ /dev/null @@ -1,145 +0,0 @@ -#ifndef _NN_KERNEL_H_ -#define _NN_KERNEL_H_ - -#include - -#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 deleted file mode 100644 index a55ac20..0000000 --- a/benchmarks/CUDA/NN/README.GPGPU-Sim +++ /dev/null @@ -1,2 +0,0 @@ -make -./gpgpu_ptx_sim__neuralnet 28 diff --git a/benchmarks/CUDA/NN/data/lw1.wei b/benchmarks/CUDA/NN/data/lw1.wei deleted file mode 100644 index d6e140e..0000000 Binary files a/benchmarks/CUDA/NN/data/lw1.wei and /dev/null differ diff --git a/benchmarks/CUDA/NN/data/lw2.wei b/benchmarks/CUDA/NN/data/lw2.wei deleted file mode 100644 index 1a340ea..0000000 Binary files a/benchmarks/CUDA/NN/data/lw2.wei and /dev/null differ diff --git a/benchmarks/CUDA/NN/data/lw3.wei b/benchmarks/CUDA/NN/data/lw3.wei deleted file mode 100644 index e86a6b6..0000000 Binary files a/benchmarks/CUDA/NN/data/lw3.wei and /dev/null differ diff --git a/benchmarks/CUDA/NN/data/lw4.wei b/benchmarks/CUDA/NN/data/lw4.wei deleted file mode 100644 index 17cf469..0000000 Binary files a/benchmarks/CUDA/NN/data/lw4.wei and /dev/null differ diff --git a/benchmarks/CUDA/NN/data/t10k-images-idx3-ubyte b/benchmarks/CUDA/NN/data/t10k-images-idx3-ubyte deleted file mode 100644 index 1170b2c..0000000 Binary files a/benchmarks/CUDA/NN/data/t10k-images-idx3-ubyte and /dev/null differ -- cgit v1.3