 			//copy over our data
 			memcpy(pres_temp, &pres[offset], pitch_host);
 			cudaErrorCheck( cudaMemcpy(dPres, pres_temp, pitch_host, cudaMemcpyHostToDevice) );
-			memcpy2D_in(mass_temp, pitch_device, mass_frac, pitch_host_real, offset, pitch_host, NSP);
-			cudaErrorCheck( cudaMemcpy2D(dMass, pitch_device, mass_frac,
-							pitch_host, pitch_host, NSP, cudaMemcpyHostToDevice) );
+			memcpy2D_in(mass_temp, padded, mass_frac, num, offset, pitch_host, NSP);
+			cudaErrorCheck( cudaMemcpy2D(dMass, pitch_device, mass_temp,
+							pitch_device, pitch_host, NSP, cudaMemcpyHostToDevice) );
 
 			if (eval_rates)
 			{
 				//eval dydt
 				//this gets us all arrays but the Jacobian
-				k_dydt<<<grid_num, TARGET_BLOCK_SIZE, SHARED_SIZE>>>(num, dPres, dMass, h_mem->dy, d_mem);
+				k_dydt<<<grid_num, TARGET_BLOCK_SIZE, SHARED_SIZE>>>(num_cond, dPres, dMass, h_mem->dy, d_mem);
 
 				check_err();
 
 				//copy back
-				cudaErrorCheck( cudaMemcpy2D(conc_temp, pitch_host, h_mem->conc, pitch_device,
+				cudaErrorCheck( cudaMemcpy2D(conc_temp, pitch_device, h_mem->conc, pitch_device,
 								pitch_host, NSP, cudaMemcpyDeviceToHost) );
-				memcpy2D_out(conc, pitch_host_real, conc_temp, pitch_host, offset, pitch_host, NSP);
-
-				cudaErrorCheck( cudaMemcpy2D(fwd_rxn_rates, pitch_host, h_mem->fwd_rates, pitch_device,
+				memcpy2D_out(conc, num, conc_temp, padded, offset, pitch_host, NSP);
+
+				cudaErrorCheck( cudaMemcpy2D(fwd_temp, pitch_device, h_mem->fwd_rates, pitch_device,
 											 pitch_host, FWD_RATES, cudaMemcpyDeviceToHost) );
-				memcpy2D_out(fwd_rxn_rates, pitch_host_real, fwd_temp, pitch_host, offset, pitch_host, FWD_RATES);
+				memcpy2D_out(fwd_rxn_rates, num, fwd_temp, padded, offset, pitch_host, FWD_RATES);
 
 				#if REV_RATES != 0
-					cudaErrorCheck( cudaMemcpy2D(rev_rxn_rates, pitch_host, h_mem->rev_rates,
+					cudaErrorCheck( cudaMemcpy2D(rev_temp, pitch_device, h_mem->rev_rates,
 											pitch_device, pitch_host, REV_RATES, cudaMemcpyDeviceToHost) );
-					memcpy2D_out(rev_rxn_rates, pitch_host_real, rev_temp, pitch_host, offset, pitch_host, REV_RATES);
+					memcpy2D_out(rev_rxn_rates, num, rev_temp, padded, offset, pitch_host, REV_RATES);
 				#endif
 
 				#if PRES_MOD_RATES != 0
-				cudaErrorCheck( cudaMemcpy2D(pres_mod, pitch_host, h_mem->pres_mod, pitch_device, pitch_host,
+				cudaErrorCheck( cudaMemcpy2D(pres_mod_temp, pitch_device, h_mem->pres_mod, pitch_device, pitch_host,
 												PRES_MOD_RATES, cudaMemcpyDeviceToHost) );
-				memcpy2D_out(pres_mod, pitch_host_real, pres_mod_temp, pitch_host, offset, pitch_host, PRES_MOD_RATES);
+				memcpy2D_out(pres_mod, num, pres_mod_temp, padded, offset, pitch_host, PRES_MOD_RATES);
 				#endif
 
-				cudaErrorCheck( cudaMemcpy2D(spec_rates, pitch_host, h_mem->spec_rates, pitch_device,
+				cudaErrorCheck( cudaMemcpy2D(spec_temp, pitch_device, h_mem->spec_rates, pitch_device,
 												pitch_host, NSP, cudaMemcpyDeviceToHost) );
-				memcpy2D_out(spec_rates, pitch_host_real, spec_temp, pitch_host, offset, pitch_host, NSP);
-
-				cudaErrorCheck( cudaMemcpy2D(dy, pitch_host, h_mem->dy, pitch_device, pitch_host,
+				memcpy2D_out(spec_rates, num, spec_temp, padded, offset, pitch_host, NSP);
+
+				cudaErrorCheck( cudaMemcpy2D(dy_temp, pitch_device, h_mem->dy, pitch_device, pitch_host,
 												NSP, cudaMemcpyDeviceToHost) );
-				memcpy2D_out(dy, pitch_host_real, dy_temp, pitch_host, offset, pitch_host, NSP);
+				memcpy2D_out(dy, num, dy_temp, padded, offset, pitch_host, NSP);
 			}
 
 			//jacobian
-			k_eval_jacob<<<grid_num, TARGET_BLOCK_SIZE, SHARED_SIZE>>>(num, dPres, dMass, h_mem->jac, d_mem);
+			k_eval_jacob<<<grid_num, TARGET_BLOCK_SIZE, SHARED_SIZE>>>(num_cond, dPres, dMass, h_mem->jac, d_mem);
 
 			check_err();
 
 			//copy back
-			cudaErrorCheck( cudaMemcpy2D(jac, pitch_host, h_mem->jac, pitch_device,
+			cudaErrorCheck( cudaMemcpy2D(jac_temp, pitch_device, h_mem->jac, pitch_device,
 											pitch_host, NSP * NSP, cudaMemcpyDeviceToHost) );
-			memcpy2D_out(jac, pitch_host_real, jac_temp, pitch_host, offset, pitch_host, NSP * NSP);
+			memcpy2D_out(jac, num, jac_temp, padded, offset, pitch_host, NSP * NSP);
 
 			offset += num_per_kernel;
 		}

			double* pres_mod, double* spec_rates, double* dy, double* jac,
			bool eval_rates)
{
	int grid_num = (int)ceil(((double)num) / ((double)TARGET_BLOCK_SIZE));
	mechanism_memory * d_mem = 0;
	mechanism_memory * h_mem = (mechanism_memory*)malloc(sizeof(mechanism_memory));
	double* dMass = 0;

	int padded = -1;
	while (padded == -1)
	{
		grid_num = (int)ceil(((double)num_per_kernel) / ((double)TARGET_BLOCK_SIZE));
		padded = initialize_gpu_memory(num_per_kernel, TARGET_BLOCK_SIZE, grid_num, &h_mem, &d_mem, &dMass, &dPres);
		if (padded == -1)
		{
			num_kernels *= 2;
			num_per_kernel /= 2;

			free_gpu_memory(&h_mem, &d_mem, &dMass, &dPres);
			free(h_mem);
			cudaErrorCheck( cudaDeviceReset() );
		}
	}
	int offset = 0;
