__global__ void kernel_convoNonSepSh_8p(unsigned char *output, int j_dim, int r)
{
  int ic, jc;
  int k = 2*r+1 ;
  float outval0=0.0, outval1=0.0, outval2=0.0, outval3=0.0, outval4=0.0, outval5=0.0, outval6=0.0, outval7=0.0 ;
  int bdimX = blockDim.x<<3 ;
  int tidX = threadIdx.x<<3 ;
    
  // absolute coordinates of packet's base point
  int j = (__umul24(blockIdx.x,blockDim.x) + threadIdx.x) << 3 ; 
  int i = __umul24( blockIdx.y, blockDim.y) + threadIdx.y ;
  int j0= __umul24(blockIdx.x,blockDim.x)<<3 ;   // block's base point
  int idx = __umul24(i,j_dim) + j ;              // absolute index
  int idrow = threadIdx.y*(bdimX+k-1) ;          // line's offset in sh mem
  
  // shared memory declaration
  extern __shared__ unsigned char roi8p[];

  // top left
  for (int p=0; p<8; p++)
  roi8p[ idrow  + tidX +p ] = tex2D(tex_img_inc, j-r+p  , i-r) ;
  
  // top right
  if ( threadIdx.x < r ) 
	{
	  roi8p[ idrow + bdimX + threadIdx.x    ] = tex2D( tex_img_inc, j0-r +bdimX+threadIdx.x  , i-r ) ;
	  roi8p[ idrow + bdimX + threadIdx.x +r ] = tex2D( tex_img_inc, j0   +bdimX+threadIdx.x  , i-r ) ;
	}
  // bottom left
  if ( threadIdx.y < k-1 )
	{
	  idrow = (threadIdx.y+blockDim.y)*(bdimX+k-1) ;
	  for (int p=0; p<8; p++)
		roi8p[ idrow + tidX +p  ] = tex2D( tex_img_inc, j-r+p  , i+blockDim.y-r ) ;
	  // bottom right
	  if ( threadIdx.x < r ) 
		{
		  roi8p[ idrow + bdimX +threadIdx.x   ] = tex2D( tex_img_inc, j0-r +bdimX +threadIdx.x, i+blockDim.y-r ) ;
		  roi8p[ idrow + bdimX +threadIdx.x +r] = tex2D( tex_img_inc, j0   +bdimX +threadIdx.x, i+blockDim.y-r ) ;
		}
	}
  __syncthreads();
 
  // computations
  for (ic=0 ; ic<k ; ic++)
	for( jc=0 ; jc<k ; jc++)
	  {
		int baseRoi = __umul24(ic+threadIdx.y,(bdimX+k-1)) + jc+tidX ;
		float valMask = mask[ __umul24(ic,k)+jc ] ;
		outval0 += valMask*roi8p[ baseRoi ] ;
		outval1 += valMask*roi8p[ baseRoi +1 ] ;
		outval2 += valMask*roi8p[ baseRoi +2 ] ;
		outval3 += valMask*roi8p[ baseRoi +3 ] ;
		outval4 += valMask*roi8p[ baseRoi +4 ] ;
		outval5 += valMask*roi8p[ baseRoi +5 ] ;
		outval6 += valMask*roi8p[ baseRoi +6 ] ;
		outval7 += valMask*roi8p[ baseRoi +7 ] ;
	  }
	
  // multiple output --> global mem
  output[ idx++ ] = outval0 ;
  output[ idx++ ] = outval1 ;
  output[ idx++ ] = outval2 ;
  output[ idx++ ] = outval3 ;
  output[ idx++ ] = outval4 ;
  output[ idx++ ] = outval5 ;
  output[ idx++ ] = outval6 ;
  output[ idx   ] = outval7 ;
}