 	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;

		//special handling for kernels that would require too much memory

		grid_num = (int)ceil(((double)num_per_kernel) / ((double)TARGET_BLOCK_SIZE));
		int pitch_host_real = num * sizeof(double);
		//first we need temporary copies on the host memory for memory transfer
		double* pres_temp = (double*)malloc(pitch_device);
		double* mass_temp = (double*)malloc(NSP * pitch_device);

			//copy over our data
			memcpy(pres_temp, &pres[offset], pitch_host);
			cudaErrorCheck( cudaMemcpy(dPres, pres_temp, pitch_host, 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_cond, dPres, dMass, h_mem->dy, d_mem);

				check_err();

				//copy back
				cudaErrorCheck( cudaMemcpy2D(conc_temp, pitch_device, h_mem->conc, pitch_device,
								pitch_host, NSP, cudaMemcpyDeviceToHost) );
				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, num, fwd_temp, padded, offset, pitch_host, FWD_RATES);

				#if REV_RATES != 0
					cudaErrorCheck( cudaMemcpy2D(rev_temp, pitch_device, h_mem->rev_rates,
											pitch_device, pitch_host, REV_RATES, cudaMemcpyDeviceToHost) );
					memcpy2D_out(rev_rxn_rates, num, rev_temp, padded, offset, pitch_host, REV_RATES);
				#endif

				#if PRES_MOD_RATES != 0
				cudaErrorCheck( cudaMemcpy2D(pres_mod_temp, pitch_device, h_mem->pres_mod, pitch_device, pitch_host,
												PRES_MOD_RATES, cudaMemcpyDeviceToHost) );
				memcpy2D_out(pres_mod, num, pres_mod_temp, padded, offset, pitch_host, PRES_MOD_RATES);
				#endif

				cudaErrorCheck( cudaMemcpy2D(spec_temp, pitch_device, h_mem->spec_rates, pitch_device,
												pitch_host, NSP, cudaMemcpyDeviceToHost) );
				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, num, dy_temp, padded, offset, pitch_host, NSP);
			}

			//jacobian
			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_temp, pitch_device, h_mem->jac, pitch_device,
											pitch_host, NSP * NSP, cudaMemcpyDeviceToHost) );
			memcpy2D_out(jac, num, jac_temp, padded, offset, pitch_host, NSP * NSP);

			offset += num_per_kernel;
		}
