 			break;
 		}
 
-		CUDA_CHECK(cudaMemcpyAsync(&PVec[0], d_P, sizeof(double) * N, cudaMemcpyDeviceToHost, stream));
-	}
+		for (auto& v : PVec) v -= 2.0;
+	}
+
+	CUDA_CHECK(cudaHostUnregister(Tpot.data()));
+	CUDA_CHECK(cudaHostUnregister(MR.data()));
+	CUDA_CHECK(cudaHostUnregister(PVec.data()));
 
 	CUDA_CHECK(cudaStreamSynchronize(stream));
 
 	// Calculate averages
 
-	auto Tpot = tp.Result();
-	auto MR = mr.Result();
+	Tpot = tp.Result();
+	MR = mr.Result();
 
 	// Copy averages to GPU for final calculation
 	CUDA_CHECK(cudaMemcpyAsync(d_Tpot, &Tpot[0], sizeof(double) * N, cudaMemcpyHostToDevice, stream));

	InitializeArray<double>(d_Tpot, himan::MissingDouble(), N, stream);
	InitializeArray<double>(d_MR, himan::MissingDouble(), N, stream);

	while (true)
	{
		auto TVec = h->VerticalValue(param("T-K"), PVec);
		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));


	if (idx < N)
	{
		const double T = d_T[idx];
		const double P = d_P[idx];
		const double RH = d_RH[idx];

		ASSERT((T > 150 && T < 350) || IsMissingDouble(T));
		ASSERT((P > 100 && P < 1500) || IsMissingDouble(P));
		ASSERT((RH >= 0 && RH < 102) || IsMissingDouble(RH));

		d_Tpot[idx] = metutil::Theta_(T, 100 * P);
		d_MR[idx] = metutil::smarttool::MixingRatio_(T, RH, 100 * P);

		d_P[idx] = P - 2.0;
	}
}

