 	CUDA_CHECK(cudaFree(d_prevMR));
 	CUDA_CHECK(cudaFree(d_range));
 	CUDA_CHECK(cudaFree(d_P500m));
+	CUDA_CHECK(cudaFree(d_Tsurface));
+	CUDA_CHECK(cudaFree(d_RHsurface));
+	CUDA_CHECK(cudaFree(d_Tfirst));
+	CUDA_CHECK(cudaFree(d_RHfirst));
+
+	float* d_T = 0;
+	float* d_TD = 0;
 
 	CUDA_CHECK(cudaMalloc((float**)&d_TD, N * sizeof(float)));
 	CUDA_CHECK(cudaMalloc((float**)&d_T, N * sizeof(float)));


	CUDA_CHECK(cudaMemcpyAsync(T.data(), d_T, sizeof(float) * N, cudaMemcpyDeviceToHost, stream));
	CUDA_CHECK(cudaMemcpyAsync(TD.data(), d_TD, sizeof(float) * N, cudaMemcpyDeviceToHost, stream));
	CUDA_CHECK(cudaStreamSynchronize(stream));

	CUDA_CHECK(cudaFree(d_meanTpot));


		ASSERT((P > 100 && P < 1500) || IsMissing(P));

		float T = Tpot * pow((P / 1000.), 0.2854);
		const float Es = metutil::Es_<float>(T);  // Saturated water vapor pressure
		const float E = metutil::E_<float>(MR, 100 * P);
		const float RH = fminf(102., E / Es * 100);

		d_TD[idx] = metutil::DewPointFromRH_<float>(T, RH);
		d_T[idx] = T;

		if (isnan(T) || isnan(d_TD[idx]))
		{
			// half is mangling the missing value-nan
			d_T[idx] = himan::MissingFloat();
			d_TD[idx] = himan::MissingFloat();
		}
	}
}

