 				sample = himan::MissingFloat();
 			}
 
-			d_sampled[idx * sampleCount + i] = himan::MissingFloat();
+			d_sampled[idx * sampleCount + i] = __float2half(himan::MissingFloat());
 
 			for (int j = 1; j < levelCount; j++)
 			{
-				const float x1 = d_x[idx * levelCount + j - 1];
-				const float y1 = d_y[idx * levelCount + j - 1];
-				const float x2 = d_x[idx * levelCount + j];
-				const float y2 = d_y[idx * levelCount + j];
+				const int ii = idx * levelCount + j - 1;
+				const float x1 = d_x0[idx] - __half2float(d_x[ii]);
+				const float y1 = d_yfirst[idx] - __half2float(d_y[ii]);
+				const float x2 = d_x0[idx] - __half2float(d_x[ii + 1]);
+				const float y2 = d_yfirst[idx] - __half2float(d_y[ii + 1]);
 
 				if (x1 >= sample && x2 <= sample)
 				{
-					d_sampled[idx * sampleCount + h] = interpolation::Linear<float>(sample, x1, x2, y1, y2);
+					const float val = interpolation::Linear<float>(sample, x1, x2, y1, y2);
+
+					if (h == 0)
+					{
+						d_first[idx] = val;
+					}
+
+					d_sampled[idx * sampleCount + h] = __float2half(d_first[idx] - val);
 					h++;
 					break;
 				}

		if (IsValid(d_Tpot[idx]))
		{
			// trapezoidal integration

			d_range[idx] += 1;
			d_meanTpot[idx] += (d_prevTpot[idx] + d_Tpot[idx]) * 0.5;
			d_meanMR[idx] += (d_prevMR[idx] + d_MR[idx]) * 0.5;

	float* d_meanTpot = 0;
	float* d_meanMR = 0;
	float* d_range = 0;

	CUDA_CHECK(cudaMalloc((float**)&d_Tpot, N * sizeof(float)));
	CUDA_CHECK(cudaMalloc((float**)&d_MR, N * sizeof(float)));
