 	InitializeArray<double>(d_Tpot, himan::MissingDouble(), N, stream);
 	InitializeArray<double>(d_MR, himan::MissingDouble(), N, stream);
 
+	std::vector<double> Tpot(N, himan::MissingDouble());
+	std::vector<double> MR(N, himan::MissingDouble());
+
+	CUDA_CHECK(cudaHostRegister(reinterpret_cast<void*>(Tpot.data()), sizeof(double) * N, 0));
+	CUDA_CHECK(cudaHostRegister(reinterpret_cast<void*>(MR.data()), sizeof(double) * N, 0));
+	CUDA_CHECK(cudaHostRegister(reinterpret_cast<void*>(PVec.data()), sizeof(double) * N, 0));
+
+	CUDA_CHECK(cudaMemcpyAsync(d_P, &PVec[0], sizeof(double) * N, cudaMemcpyHostToDevice, stream));
+
 	while (true)
 	{
 		auto TVec = h->VerticalValue(param("T-K"), PVec);
+		CUDA_CHECK(cudaMemcpyAsync(d_T, &TVec[0], sizeof(double) * N, cudaMemcpyHostToDevice, stream));
+
 		auto RHVec = h->VerticalValue(param("RH-PRCNT"), PVec);
-
-		CUDA_CHECK(cudaMemcpyAsync(d_T, &TVec[0], sizeof(double) * N, cudaMemcpyHostToDevice, stream));
 		CUDA_CHECK(cudaMemcpyAsync(d_RH, &RHVec[0], sizeof(double) * N, cudaMemcpyHostToDevice, stream));
-		CUDA_CHECK(cudaMemcpyAsync(d_P, &PVec[0], sizeof(double) * N, cudaMemcpyHostToDevice, stream));
 
 		MixingRatioKernel<<<gridSize, blockSize, 0, stream>>>(d_T, d_P, d_RH, d_Tpot, d_MR, N);
-
-		std::vector<double> Tpot(N, himan::MissingDouble());
-		std::vector<double> MR(N, himan::MissingDouble());
-
-		CUDA_CHECK(cudaStreamSynchronize(stream));
 
 		CUDA_CHECK(cudaMemcpyAsync(&Tpot[0], d_Tpot, sizeof(double) * N, cudaMemcpyDeviceToHost, stream));
 		CUDA_CHECK(cudaMemcpyAsync(&MR[0], d_MR, sizeof(double) * N, cudaMemcpyDeviceToHost, stream));

			break;
		}

		CUDA_CHECK(cudaMemcpyAsync(&PVec[0], d_P, sizeof(double) * N, cudaMemcpyDeviceToHost, stream));
	}

	CUDA_CHECK(cudaStreamSynchronize(stream));

	// Calculate averages

	auto Tpot = tp.Result();
	auto MR = mr.Result();

	// Copy averages to GPU for final calculation
	CUDA_CHECK(cudaMemcpyAsync(d_Tpot, &Tpot[0], sizeof(double) * N, cudaMemcpyHostToDevice, stream));

		ASSERT(d_T[idx] > 100);
		ASSERT(d_T[idx] < 350 || IsMissingDouble(d_T[idx]));

		double T = metutil::LiftLCLA_(d_P[idx] * 100, d_T[idx], d_PLCL[idx] * 100, d_Ptarget.values[idx] * 100);

		ASSERT(T > 100);
		ASSERT(T < 350 || IsMissingDouble(T));
