  * \param[inout]	temp_red		temperatures of red cells
  * \param[out]		bl_norm_L2	array with residual information for blocks
  */
+#ifdef TEXTURE
+__global__ void red_kernel (const Real * temp_black, Real * temp_red, Real * bl_norm_L2)
+#else
 __global__ void red_kernel (const Real * aP, const Real * aW, const Real * aE,
 														const Real * aS, const Real * aN, const Real * b,
 														const Real * temp_black, Real * temp_red,
 														Real * bl_norm_L2)
+#endif
 {	
 	uint row = 1 + (blockIdx.y * blockDim.y) + threadIdx.y;
 	uint col = 1 + (blockIdx.x * blockDim.x) + threadIdx.x;

		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));
		

		}
		
		// update black cells
		#ifdef TEXTURE
		black_kernel <<<dimGrid, dimBlock>>> (temp_red_d, temp_black_d, bl_norm_L2_d);
		#else
		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);
		#endif
		
		// sync threads (needed?)
		//CUDA_SAFE_CALL (cudaThreadSynchronize());
