 	uint ind_black = col * ((NUM >> 1) + 2) + row;  					// local (black) index
 	uint ind = 2 * row - ((col + 1) & 1) - 1 + NUM * (col - 1);	// global index
 	
-	Real res = b[ind] + (aW[ind] * temp_red[row + (col - 1) * ((NUM >> 1) + 2)]
-									   + aE[ind] * temp_red[row + (col + 1) * ((NUM >> 1) + 2)]
-									   + aS[ind] * temp_red[row - ((col + 1) & 1) + col * ((NUM >> 1) + 2)]
-									   + aN[ind] * temp_red[row + (col & 1) + col * ((NUM >> 1) + 2)]);
-	
 	Real temp_old = temp_black[ind_black];
-	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_red[row + (col - 1) * ((NUM >> 1) + 2)]
+				    + get_tex(aE_t, ind) * temp_red[row + (col + 1) * ((NUM >> 1) + 2)]
+				    + get_tex(aS_t, ind) * temp_red[row - ((col + 1) & 1) + col * ((NUM >> 1) + 2)]
+				    + get_tex(aN_t, ind) * temp_red[row + (col & 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_red[row + (col - 1) * ((NUM >> 1) + 2)]
+				    + aE_t[ind] * temp_red[row + (col + 1) * ((NUM >> 1) + 2)]
+				    + aS_t[ind] * temp_red[row - ((col + 1) & 1) + col * ((NUM >> 1) + 2)]
+				    + aN_t[ind] * temp_red[row + (col & 1) + col * ((NUM >> 1) + 2)]);
+	
+	Real temp_new = temp_old * (1.0 - omega) + omega * (res / aP_t[ind]);
+	#endif
 	
 	temp_black[ind_black] = temp_new;
 	res = temp_old - temp_new;

		}
		
		// update black cells
		black_kernel <<<dimGrid, dimBlock>>> (aP_d, aW_d, aE_d, aS_d, aN_d, b_d, temp_red_d, temp_black_d, bl_norm_L2_d);
		
		// sync threads (needed?)
		//CUDA_SAFE_CALL (cudaThreadSynchronize());

#include <cutil.h>

/** Double precision */
#define DOUBLE

#ifdef DOUBLE
	#define Real double
#else
	#define Real float
#endif

/** Use texture memory */
#define TEXTURE

typedef unsigned int uint;

