 
 	if (idx < N)
 	{
-		double T = d_T[idx];
-		double P = d_P[idx];
-		double RH = d_RH[idx];
+		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));
 
-		if (IsMissingDouble(T) || IsMissingDouble(P) || IsMissingDouble(RH))
-		{
-			d_P[idx] = MissingDouble();
-		}
-		else
-		{
-			d_Tpot[idx] = metutil::Theta_(T, 100 * P);
-			d_MR[idx] = metutil::smarttool::MixingRatio_(T, RH, 100 * P);
-
-			d_P[idx] = P - 2.0;
-		}
+		d_Tpot[idx] = metutil::Theta_(T, 100 * P);
+		d_MR[idx] = metutil::smarttool::MixingRatio_(T, RH, 100 * P);
+
+		d_P[idx] = P - 2.0;
 	}
 }
 

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

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

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

			break;
		}

		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

	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));
