 	case 'g':
 		if (var == 'P')
 		{
-			result_d_ro = (def.ro0_g / def.P_atm) * (T_0 / T);
+			result_d_ro = (gpu_def->ro0_g / gpu_def->P_atm) * (T_0 / T);
 		} 
 		else if (var == 'T')
 		{
-			result_d_ro = def.ro0_g * (P / def.P_atm) * (-1) * (T_0 / T) / T;
+			result_d_ro = gpu_def->ro0_g * (P / gpu_def->P_atm) * (-1) * (T_0 / T) / T;
 		}
 		else 
 		{

	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)))
	{
		int local1 = set_boundary_basic_coordinate(i, j, k, def);
		int local = i + j * (def.locNx) + k * (def.locNx) * (def.locNy);

		if (j == 0)
		{
			HostArraysPtr.T[local] = 400;
		}
		else if(j == (def.locNy) - 1)
		{
			HostArraysPtr.T[local] = 273;
		}
		else
		{
			// Будем считать границы области не теплопроводящими
			HostArraysPtr.T[local] = HostArraysPtr.T[local1];
		}

		test_positive(HostArraysPtr.T[local], __FILE__, __LINE__);
	}
}


			dF[3] = 0.;
			dF[4] = 0.;

			dF[0 + n] = ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'w');
			dF[1 + n] = 0.;
			dF[2 + n] = 0.;
			dF[3 + n] = d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'w', 'P') * DevArraysPtr.S_w[local];
			dF[4 + n] = d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'w', 'T') * DevArraysPtr.S_w[local];

			dF[0 + n * 2] = 0.;
			dF[1 + n * 2] = ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'n');
			dF[2 + n * 2] = 0.;
			dF[3 + n * 2] = d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'n', 'P') * DevArraysPtr.S_n[local];
			dF[4 + n * 2] = d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'n', 'T') * DevArraysPtr.S_n[local];

			dF[0 + n * 3] = 0.;
			dF[1 + n * 3] = 0.;
			dF[2 + n * 3] = ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'g');
			dF[3 + n * 3] = d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'g', 'P') * DevArraysPtr.S_g[local];
			dF[4 + n * 3] = d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'g', 'T') * DevArraysPtr.S_g[local];

			dF[0 + n * 4] = DevArraysPtr.m[local] * (ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'w') * DevArraysPtr.H_w[local] - DevArraysPtr.P_w[local]);
			dF[1 + n * 4] = DevArraysPtr.m[local] * (ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'n') * DevArraysPtr.H_n[local] - DevArraysPtr.P_w[local]);
			dF[2 + n * 4] = DevArraysPtr.m[local] * (ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'g') * DevArraysPtr.H_g[local] - DevArraysPtr.P_w[local]);
			dF[3 + n * 4] = DevArraysPtr.m[local] * (DevArraysPtr.S_w[local] * (d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'w', 'P') * DevArraysPtr.H_w[local] - 1.) 
				+ DevArraysPtr.S_n[local] * (d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'n', 'P') * DevArraysPtr.H_n[local] - 1.)  
				+ DevArraysPtr.S_g[local] * (d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'g', 'P') * DevArraysPtr.H_g[local] - 1.))
				+ (1. - DevArraysPtr.m[local]) * (-1);
			dF[4 + n * 4] = DevArraysPtr.m[local] * (DevArraysPtr.S_w[local] * (d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'w', 'T') * DevArraysPtr.H_w[local] 
				+ ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'w') * c_w(DevArraysPtr.T[local])) 
				+ DevArraysPtr.S_n[local] * (d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'n', 'T') * DevArraysPtr.H_n[local]
				+ ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'n') * c_n(DevArraysPtr.T[local]))
				+ DevArraysPtr.S_g[local] * (d_ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'g', 'T') * DevArraysPtr.H_g[local]
				+ ro(DevArraysPtr.P_w[local], DevArraysPtr.T[local], 'g') * c_g(DevArraysPtr.T[local])))
				+ (1. - DevArraysPtr.m[local]) * ro_r * c_r(DevArraysPtr.T[local]);

			reverse_matrix(dF, n);
			mult_matrix_vector(correction, dF, F, n);

			DevArraysPtr.S_w[local] = DevArraysPtr.S_w[local] - correction[0];
			DevArraysPtr.S_n[local] = DevArraysPtr.S_n[local] - correction[1];
			DevArraysPtr.S_g[local] = DevArraysPtr.S_g[local] - correction[2];
			DevArraysPtr.P_w[local] = DevArraysPtr.P_w[local] - correction[3];
			DevArraysPtr.T[local] = DevArraysPtr.T[local] - correction[4];
			assign_H(DevArraysPtr, local);
		}

		// Обновление значения суммарной энергии, т.к. оно больше не понадобится
		// !!! Лучше вынести в отдельную функцию (просто обменять указатели).
		DevArraysPtr.E[local] = DevArraysPtr.E_new[local];

		delete[] F;
		delete[] correction;
		delete[] dF;

		device_test_S(DevArraysPtr.S_w[local], __FILE__, __LINE__);
		device_test_S(DevArraysPtr.S_n[local], __FILE__, __LINE__);
		device_test_positive(DevArraysPtr.P_w[local], __FILE__, __LINE__);
		device_test_positive(DevArraysPtr.T[local], __FILE__, __LINE__);
		device_test_nan(DevArraysPtr.E[local], __FILE__, __LINE__);
	}
}
