 	memset(E, 0.0, sizeof(float) * height_A * width_B);
 	memset(F, 0.0, sizeof(float) * height_A * width_B);
 
+	Type m_Mode;
 
 	//
 	srand((unsigned)time(0));

bool CheckAnswer(const float* _C, const float* _D, unsigned int size)
{
	bool isRight = true;
	for (int i = 0; i < size && isRight == true; ++i)
	{
		if (_C[i] != _D[i])
			isRight = false;
	}

	return isRight;

	MatrixMulCPU(C, A, B, width_A, height_A, width_B, height_B);
	printf("CPU use time : %dms\n", GetTickCount() - tick1);


		//GPU 
		printf("GPU  normal matrix multiplication...\n");
		m_Mode = Mode1;

		cudaError_t cudaStatus = addWithCuda(D, A, B, width_A, height_A, width_B, height_B, m_Mode);
		if (cudaStatus != cudaSuccess)
		{
			fprintf(stderr, "addWithCuda failed!\n");
			return 1;
		}

		printf("GPU  with shared memory and matrix being blocked matrix multiplication...\n");
		m_Mode = Mode2;

		cudaStatus = addWithCuda(E, A, B, width_A, height_A, width_B, height_B, m_Mode);
		if (cudaStatus != cudaSuccess)
		{
			fprintf(stderr, "addWithCuda failed!\n");
			return 1;
		}
	
		//GPU mode3 with texture memory

		cudaChannelFormatDesc channelDescA = cudaCreateChannelDesc((int)sizeof(float) * 8, 0, 0, 0, cudaChannelFormatKindFloat);
		cudaChannelFormatDesc channelDescB = cudaCreateChannelDesc((int)sizeof(float) * 8, 0, 0, 0, cudaChannelFormatKindFloat);

		cudaArray* mat_A;
		cudaArray* mat_B;

		cudaMallocArray(&mat_A, &channelDescA, width_A, height_A);
		cudaMallocArray(&mat_B, &channelDescB, width_B, height_B);

		cudaMemcpyToArray(mat_A, 0, 0, A, sizeof(float) * height_A * width_A, cudaMemcpyHostToDevice);
		cudaMemcpyToArray(mat_B, 0, 0, B, sizeof(float) * height_B * width_B, cudaMemcpyHostToDevice);

		//texA.addressMode[0] = cudaAddressModeWrap;
		//texA.addressMode[1] = cudaAddressModeWrap;
		texA.filterMode = cudaFilterModePoint;
		texA.normalized = false;
		//texB.addressMode[0] = cudaAddressModeWrap;
		//texB.addressMode[1] = cudaAddressModeWrap;
		texB.filterMode = cudaFilterModePoint;
		texB.normalized = false;

		cudaBindTextureToArray(texA, mat_A, channelDescA);
		cudaBindTextureToArray(texB, mat_B, channelDescB);

		float* d_C = NULL;
		cudaMalloc(&d_C, width_B * height_A * sizeof(float));

		int block_size = 32;

		dim3 Threads(block_size, block_size);
		dim3 Blocks(width_B / block_size, height_A / block_size);

		cudaEvent_t start3, stop3;
		cudaEventCreate(&start3);
		cudaEventCreate(&stop3);
		cudaEventRecord(start3, 0);

		MatrixMul << < Threads, Blocks >> >(d_C, width_B, height_A);

		cudaEventRecord(stop3, 0);
		cudaEventSynchronize(stop3);

		float elaspsedTime3;
		cudaEventElapsedTime(&elaspsedTime3, start3, stop3);
		printf("GPU with Texutre Memory time : %3.1fms \n", elaspsedTime3);
		cudaMemcpy(F, d_C, sizeof(float) * width_B * height_A, cudaMemcpyDeviceToHost);

		cudaFree(d_C);
		cudaFree(mat_A);
		cudaFree(mat_B);

	
	//GPU, CPU ĽǷͬ
	printf("Checking answer...\n");
	if (!CheckAnswer(E, D, height_A * width_B))
		printf("The answer1 is wrong!");
	else printf("The answer1 is right!");

	if (!CheckAnswer(E, F, height_A * width_B))
		printf("The answer2 is wrong!");
	else printf("The answer2 is right!");

	//if (!CheckAnswer(C, F, height_A * width_B))
	//	printf("The answer3 is wrong!");
	//else printf("The answer3 is right!");

	// cudaDeviceReset must be called before exiting in order for profiling and
	// tracing tools such as Nsight and Visual Profiler to show complete traces.
