 	return;
 }
 
+void initEigenvectors(float* container) {
+	FILE* ptr=fopen("eigenvectors.csv", "r");
+	char* buffer=(char*)malloc(sizeof(char));
+	for(int j=0; j<DataLen; ++j) {
+		for(int i=0; i<Reduced_Data_Length; ++i) {
+			fscanf(ptr, "%f", &container[j*Reduced_Data_Length+i]);
+			fscanf(ptr, "%c", buffer);
+		}
+	}
+	free(buffer);
+	return;
+}
+
+
 void initPCAKNN() {
+
 	digits=(int*)malloc(TrainDataNum*sizeof(int));
 	initDigits(digits);
+	cudaMalloc((void **) &digitDevice, digits, TrainDataNum*sizeof(int), cudaMemcpyHostToDevice);
+	cuCheck(__LINE__);
 	int distantPCASize=sizeof(float)*TrainDataNum*15;
 	distantPCAHost=(float*)malloc(distantPCASize);
 	cudaMalloc((void**) &distantPCAKernel, distantPCASize);
 	cuCheck(__LINE__);
-	//init pcatraindata
+	initEigenvectors(trainDataPCAHost);
+	cudaMalloc((void**) &trainDataPCAKernel, trainDataSize);
+	cuCheck(__LINE__);
+	cudaMemcpy(trainDataKernel, trainDataHost, trainDataSize, cudaMemcpyHostToDevice);
+	cuCheck(__LINE__);
 	//init dist
 }
 
 
-__global__ void knn(float* trainDataKernel, float* distantKernel, int count) {
-	__shared__ float digit[DataWidth][DataWidth];
-	__shared__ float train[DataWidth][DataWidth];
+__global__ void knnPCA(float* trainDataPCAKernel, float* distantPCAKernel, int count) {
+	__shared__ float digit[Reduced_Data_Length];
+	__shared__ float train[Reduced_Data_Length];
 	int tx=threadIdx.x;
-	int ty=threadIdx.y;
 	int bx=blockIdx.x;
-
-	train[tx][ty]=trainDataKernel[bx*DataLen+ty*DataWidth+tx];
-
+	train[tx]=trainDataPCAKernel[bx*Reduced_Data_Length+tx];
 	for(int i=0; i<count; ++i) {
-		digit[tx][ty]=testDigit[i*DataLen+ty*DataWidth+tx];
+		digit[tx]=testPCADigit[i*Reduced_Data_Length+tx];
 		__syncthreads();
-
-		float cur=digit[tx][ty]-train[tx][ty];
+		float cur=digit[tx]-train[tx];
 		cur=cur*cur;
-		digit[tx][ty]=cur;
-		for(int stride=16; stride>0; stride/=2) {
+		digit[tx]=cur;
+		for(int stride=32; stride>0; stride/=2) {
 			__syncthreads();
-			if(tx<stride&&ty<stride) {
-				digit[tx][ty]+=digit[tx+stride][ty]+digit[tx+stride][ty+stride]+digit[tx][ty+stride];
+			if(tx<stride) {
+				digit[tx]=digit[tx]+digit[tx+stride];
 			}
 		}
 		__syncthreads();
 		if(tx==0) {
-			distantKernel[TrainDataNum*i+bx]=int(digit[0][0])+float(bx)/10000;
-		}
-
+			distantPCAKernel[Reduced_Data_Length*i+bx]=digit;
+		}
 		__syncthreads();
 	}
 }
+
+
+// __global__ void knn(float* trainDataKernel, float* distantKernel, int count) {
+// 	__shared__ float digit[DataWidth][DataWidth];
+// 	__shared__ float train[DataWidth][DataWidth];
+// 	int tx=threadIdx.x;
+// 	int ty=threadIdx.y;
+// 	int bx=blockIdx.x;
+
+// 	train[tx][ty]=trainDataKernel[bx*DataLen+ty*DataWidth+tx];
+
+// 	for(int i=0; i<count; ++i) {
+// 		digit[tx][ty]=testDigit[i*DataLen+ty*DataWidth+tx];
+// 		__syncthreads();
+
+// 		float cur=digit[tx][ty]-train[tx][ty];
+// 		cur=cur*cur;
+// 		digit[tx][ty]=cur;
+// 		for(int stride=16; stride>0; stride/=2) {
+// 			__syncthreads();
+// 			if(tx<stride&&ty<stride) {
+// 				digit[tx][ty]+=digit[tx+stride][ty]+digit[tx+stride][ty+stride]+digit[tx][ty+stride];
+// 			}
+// 		}
+// 		__syncthreads();
+// 		if(tx==0) {
+// 			distantKernel[TrainDataNum*i+bx]=int(digit[0][0])+float(bx)/10000;
+// 		}
+
+// 		__syncthreads();
+// 	}
+// }
 
 
 // input: the pointer to hold the training datas


//knn code below{{{{{{{{{{{{{{{{{{{{{{{{
int* digits=NULL;
float* trainDataHost=NULL;
float* trainDataKernel=NULL;
float* distantHost=NULL;



void setPCAConstant(float* digitHost, int count) {
	int digitSize=sizeof(float)*count*DataLen;
	int digitPCADSize=sizeof(float)*count*Reduced_Data_Length;
	float* digitDevice;
	float* digitPCADevice;
