 	uint ind_red = col * ((NUM >> 1) + 2) + row;  		// local (red) index
 	uint ind = 2 * row - (col & 1) - 1 + NUM * (col - 1);	// global index
 	
-	Real res = b[ind] + (aW[ind] * temp_black[row + (col - 1) * ((NUM >> 1) + 2)]
-									   + aE[ind] * temp_black[row + (col + 1) * ((NUM >> 1) + 2)]
-									   + aS[ind] * temp_black[row - (col & 1) + col * ((NUM >> 1) + 2)]
-									   + aN[ind] * temp_black[row + ((col + 1) & 1) + col * ((NUM >> 1) + 2)]);
-	
 	Real temp_old = temp_red[ind_red];
-	Real temp_new = temp_old * (1.0 - omega) + omega * (res / aP[ind]);
+	
+	#ifdef TEXTURE
+	Real res = get_tex(b_t, ind) 
+					 + (get_tex(aW_t, ind) * temp_black[row + (col - 1) * ((NUM >> 1) + 2)]
+				    + get_tex(aE_t, ind) * temp_black[row + (col + 1) * ((NUM >> 1) + 2)]
+				    + get_tex(aS_t, ind) * temp_black[row - (col & 1) + col * ((NUM >> 1) + 2)]
+				    + get_tex(aN_t, ind) * temp_black[row + ((col + 1) & 1) + col * ((NUM >> 1) + 2)]);
+	
+	Real temp_new = temp_old * (1.0 - omega) + omega * (res / get_tex(aP_t, ind));
+	#else
+	Real res = b_t[ind]
+					 + (aW_t[ind] * temp_black[row + (col - 1) * ((NUM >> 1) + 2)]
+				    + aE_t[ind] * temp_black[row + (col + 1) * ((NUM >> 1) + 2)]
+				    + aS_t[ind] * temp_black[row - (col & 1) + col * ((NUM >> 1) + 2)]
+				    + aN_t[ind] * temp_black[row + ((col + 1) & 1) + col * ((NUM >> 1) + 2)]);
+	
+	Real temp_new = temp_old * (1.0 - omega) + omega * (res / aP_t[ind]);
+	#endif
 	
 	temp_red[ind_red] = temp_new;
 	res = temp_old - temp_new;

		Real norm_L2 = 0.0;
		
		// update red cells
		red_kernel <<<dimGrid, dimBlock>>> (aP_d, aW_d, aE_d, aS_d, aN_d, b_d, temp_black_d, temp_red_d, bl_norm_L2_d);
		
		CUDA_SAFE_CALL (cudaMemcpy (bl_norm_L2, bl_norm_L2_d, dimGrid.x * dimGrid.y * sizeof(Real), cudaMemcpyDeviceToHost));
		

	CUDA_SAFE_CALL (cudaFree(temp_black_d));
	
	CUDA_SAFE_CALL (cudaFree(bl_norm_L2_d));
	
	#ifdef TEXTURE
	// unbind textures
	CUDA_SAFE_CALL (cudaUnbindTexture(aP_t));
	CUDA_SAFE_CALL (cudaUnbindTexture(aW_t));
	CUDA_SAFE_CALL (cudaUnbindTexture(aE_t));
	CUDA_SAFE_CALL (cudaUnbindTexture(aS_t));
	CUDA_SAFE_CALL (cudaUnbindTexture(aN_t));
	CUDA_SAFE_CALL (cudaUnbindTexture(b_t));
	#endif
	
	/////////////////////////////////
	// end timer
