 
 	if(PCA) {
 		setPCAConstant(digitHost, count);
-		getPCAAns(ans, count);
 		recognizePCA(ans, count);
 	}else {
 		cudaMemcpyToSymbol(testDigit, digitHost, testDigitSize*15);



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

	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__);
	initEigenvectors(trainDataPCAHost);
	cudaMalloc((void**) &trainDataPCAKernel, trainDataSize);
	cuCheck(__LINE__);
	cudaMemcpy(trainDataKernel, trainDataHost, trainDataSize, cudaMemcpyHostToDevice);
	cuCheck(__LINE__);
	//init dist
}


__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 bx=blockIdx.x;
	train[tx]=trainDataPCAKernel[bx*Reduced_Data_Length+tx];
	for(int i=0; i<count; ++i) {
		digit[tx]=testPCADigit[i*Reduced_Data_Length+tx];
		__syncthreads();
		float cur=digit[tx]-train[tx];
		cur=cur*cur;
		digit[tx]=cur;
		for(int stride=32; stride>0; stride/=2) {
			__syncthreads();
			if(tx<stride) {
				digit[tx]=digit[tx]+digit[tx+stride];
			}
		}
		__syncthreads();
		if(tx==0) {
			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
