 	int j = threadIdx.y + blockIdx.y * blockDim.y;
 	int k = threadIdx.z + blockIdx.z * blockDim.z;
 	
-	if ((((i != 0) && (i != (def.locNx) - 1)) || ((def.locNx) < 2)) && (j != 0) && (j != (def.locNy) - 1) && (((k != 0) && (k != (def.locNz) - 1)) || ((def.locNz) < 2)))
+	if ((((i != 0) && (i != (gpu_def->locNx) - 1)) || ((gpu_def->locNx) < 2)) && (j != 0) && (j != (gpu_def->locNy) - 1) && (((k != 0) && (k != (gpu_def->locNz) - 1)) || ((gpu_def->locNz) < 2)))
 	{
 		double Q_hw = 0, Q_hr = 0; // источниковые члены
 
-		int local=i + j * (def.locNx) + k * (def.locNx) * (def.locNy);
-
-		HostArraysPtr.E_new[local] = HostArraysPtr.E[local] + (def.dt) * (assign_T_flow(HostArraysPtr, i, j, k, def) + Q_hw + Q_hr - assign_E_flow(HostArraysPtr, i, j, k, def));
-
-		test_nan(HostArraysPtr.E_new[local], __FILE__, __LINE__);
+		int local=i + j * (gpu_def->locNx) + k * (gpu_def->locNx) * (gpu_def->locNy);
+
+		DevArraysPtr.E_new[local] = DevArraysPtr.E[local] + (gpu_def->dt) * (assign_T_flow(DevArraysPtr, i, j, k) + Q_hw + Q_hr - assign_E_flow(DevArraysPtr, i, j, k));
+
+		device_test_nan(DevArraysPtr.E_new[local], __FILE__, __LINE__);
 	}
 }
 


	h_temp = (T - T_0) / N_temp;
	
	integral += (def.P_atm / def.ro0_w);
	integral += с_w(T_0, def);
	integral += с_w(T, def);

	for(int i = 2; i < N_temp; i+=2)
		sum += с_w(T_0 + i * h_temp, def);

	sum *= 2;
	integral += sum;
	sum = 0;

	for(int i = 1; i < N_temp; i+=2)
		sum += с_w(T_0 + i * h_temp, def);

	sum *= 4;
	integral += sum;

	case 'w':
		if (var == 'P')
		{
			result_d_ro = gpu_def->ro0_w * (gpu_def->beta_w);
		} 
		else if (var == 'T')
		{
			result_d_ro = (-1) * gpu_def->ro0_w * alfa_w;
		}
		else 
		{
