 void getAns(int* ans, int count) {
 	recognize(ans, count);
 	return;
+}
+
+
+__global__ void matrixMultiply_device(float *A, float *B, float *C, int numARows, int numAColumns, int numBRows, int numBColumns, int numCRows, int numCColumns) {
+	__shared__ float sharedA[32][32];
+	__shared__ float sharedB[32][32];
+	int Row=blockIdx.y*32+threadIdx.y;
+	int Column=blockIdx.x*32+threadIdx.x;
+	float resultC=0.0f;
+	for (int i=0; i<ceil(numAColumns/32.0f); i++)
+	{
+		if (Row<numARows && i*32+threadIdx.x<numAColumns)
+			sharedA[threadIdx.y][threadIdx.x]=A[Row*numAColumns+i*32+threadIdx.x];
+		else
+			sharedA[threadIdx.y][threadIdx.x]=0.0f;
+		if (i*32+threadIdx.y<numBRows && Column<numBColumns)
+			sharedB[threadIdx.y][threadIdx.x]=B[(i*32+threadIdx.y)*numBColumns+Column];
+		else
+			sharedB[threadIdx.y][threadIdx.x]=0.0f;
+		__syncthreads();
+		for (int j=0; j<32; j++)
+		{
+			resultC+=sharedA[threadIdx.y][j]*sharedB[j][threadIdx.x];
+		}
+		__syncthreads();
+	}
+	if (Row<numCRows && Column<numCColumns)
+		C[Row*numCColumns+Column]=resultC;
+	return;
+}
+
+
+void setPCAConstant(float* digitHost, int count) {
+	int digitSize=sizeof(float)*count*DataLen);
+	int digitPCADSize=sizeof(float)*count*Reduced_Data_Length;
+	float* digitDevice;
+	float* digitPCADevice;
+	cudaMalloc((void **) &digitDevice, digitSize);
+	cudaMalloc((void **) &digitPCADevice, digitPCADSize);
+	cudaMemcpy(digitDevice, digitHost, digitSize, cudaMemcpyHostToDevice);
+	cuCheck(__LINE__);
+
+	dim3 DimGrid(ceil(Reduced_Data_Length/32.0f),ceil(count/32.0f),1);
+	dim3 DimBlock(32,32,1);
+	matrixMultiply_device<<<dimGrid, dimBlock>>>(digitDevice, EigenVectors, digitPCADevice, count, DataLen, DataLen, Reduced_Data_Length, count, Reduced_Data_Length);
+	cudaDeviceSynchronize();
+	cuCheck(__LINE__);
+
+	cudaMemcpyToSymbol(testPCADigit, digitPCADevice, Reduced_Data_Length*count);
+	cuCheck(__LINE__);
+	cudaFree(digitDevice);
+	cudaFree(digitPCADevice);
 }
 
 

#define RED_THRESHOLD 176
#define GREEN_THRESHOLD 176
#define BLUE_THRESHOLD 127
#define TrainDigits "trainData/foo.csv"
#define TrainLabels "trainData/bar.csv"


void cuCheck(int line) {
    cudaError_t err = cudaGetLastError();

		}

	}

	if(PCA) {
		setPCAConstant(digitHost, count);
		getPCAAns(ans, count);
		recognizePCA(ans, count);
	}else {
		cudaMemcpyToSymbol(testDigit, digitHost, testDigitSize*15);
		cuCheck(__LINE__);
		getAns(ans, count);
	}
	free(digitHost);
	free(checkMonoHost);	
	return ans;
