levelines_kernels.cu

   1
   2 // chemins des lignes de niveaux
   3 // longueur = 4 pixels
   4 // une ligne = un chemin
   5
   6 __constant__ int pathDi[PSIZE_I][PSIZE_J-1] =
   7   {
   8     // Q1
   9         {  -1, -1, -1},       //
  10         {  -1, -1, -1},       //
  11         {  -1, -1, -1},       //
  12         {  -1, -1, -1},       //
  13         {  -1,  0,  1},       //
  14         {   0, -1,  0},
  15         // Q4
  16         {   0,  0,  0},       //
  17         {   0,  1,  1},       //
  18         {   1,  0,  1},       //
  19         {   1,  1,  1},       //
  20         {   1,  1,  1},       //
  21         {   1,  1,  1},
  22         // Q3
  23         {   1,  1,  1},       //
  24         {   1,  1,  1},       //
  25         {   1,  1,  1},       //
  26         {   1,  1,  1},       //
  27         {   1,  0, -1},       //
  28         {   0,  1,  0},
  29         // Q2
  30         {   0,  0,  0},       //
  31         {   0, -1,  0},       //
  32         {  -1,  0, -1},       //
  33         {  -1, -1, -1},       //
  34         {  -1, -1, -1},       //
  35         {  -1, -1, -1}
  36   } ;     //
  37
  38 __constant__ int pathDj[PSIZE_I][PSIZE_J-1] =
  39   {
  40         // Q1
  41         {  0,  0,  0},       //
  42         {  0,  1,  0},
  43         {  1,  0,  1},
  44         {  1,  1,  1},
  45         {  1,  1,  1},
  46         {  1,  1,  1},
  47         // Q4
  48         {  1,  1,  1},       //
  49         {  1,  1,  1},
  50         {  1,  1,  1},
  51         {  1,  1,  1},
  52         {  1,  0, -1},
  53         {  0,  1,  0},
  54         // Q3
  55         {  0,  0,  0},       //
  56         {  0, -1,  0},
  57         { -1,  0, -1},
  58         { -1, -1, -1},
  59         { -1, -1, -1},
  60         { -1, -1, -1},
  61         // Q2
  62         { -1, -1, -1},       //
  63         { -1, -1, -1},
  64         { -1, -1, -1},
  65         { -1, -1, -1},
  66         { -1,  0,  1},
  67         {  0, -1,  0}
  68   } ;
  69
  70
  71 // declare texture reference for 2D int texture
  72 texture<int, 2, cudaReadModeElementType> tex_img_in ;
  73 texture<int, 2, cudaReadModeElementType> tex_img_estim ;
  74 texture<int, 2, cudaReadModeElementType> tex_img_lniv ;
  75
  76
  77 __global__ void kernel_init_estim_from_img_in(unsigned int * d_estim, unsigned int L, unsigned int H, unsigned int r){
  78   // coordonnes du point dans l'image
  79   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
  80   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;
  81   unsigned int pos = i*L +j ;
  82   unsigned int ic , jc, ng ;
  83
  84   if( (i>r)&&(i<H-r)&&(j>r)&&(j<L-r) ){
  85         ng = 0 ;
  86         for (ic = i - r ; ic <= i + r ; ic++){
  87           for (jc = j - r ; jc <= j + r ; jc++){
  88                 ng += tex2D(tex_img_in, jc, ic ) ;
  89           }
  90         }
  91         d_estim[ pos ] = ng/((2*r+1)*(2*r+1)) ;
  92   }
  93   /*
  94   else
  95   d_estim[ pos ] = tex2D(tex_img_in, j, i) ;
  96   */
  97 }
  98
  99 __global__ void kernel_init_estim_from_img_in_global_mem(unsigned int * d_data, unsigned int * d_estim,
 100                                                                                                                  unsigned int L, unsigned int H, unsigned int r){
 101   // coordonnes du point dans l'image
 102   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
 103   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;
 104   unsigned int pos = i*L +j ;
 105   unsigned int ic , jc, ng ;
 106
 107   if( (i>r)&&(i<H-r)&&(j>r)&&(j<L-r) ){
 108         ng = 0 ;
 109         for (ic = i - r ; ic <= i + r ; ic++){
 110           for (jc = j - r ; jc <= j + r ; jc++){
 111                 ng += d_data[ ic*L + jc ] ;
 112           }
 113         }
 114         d_estim[ pos ] = ng/((2*r+1)*(2*r+1)) ;
 115   }
 116 }
 117
 118 __global__ void kernel_estim_next_step_global_mem(unsigned int * d_estim, unsigned int * d_lniv, unsigned int L, unsigned int H, unsigned int p){
 119   // coordonnes du point dans l'image
 120   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
 121   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;
 122   unsigned int pos = i*L +j ;
 123
 124   d_estim[ pos ] = ( tex2D(tex_img_in, j, i) + p*d_lniv[ pos ] )/(1+p) ;
 125
 126 }
 127
 128 __global__ void kernel_estim_next_step_texture(unsigned int * d_estim, unsigned int L, unsigned int H, unsigned int p){
 129   // coordonnes du point dans l'image
 130   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
 131   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;
 132   unsigned int pos = i*L +j ;
 133
 134   d_estim[ pos ] = ( tex2D(tex_img_in, j, i) + p*tex2D(tex_img_lniv, j, i ))/(1+p) ;
 135
 136 }
 137
 138 __global__ void kernel_estim_next_step_global_mem(unsigned int * d_estim, unsigned int * d_lniv, unsigned int * d_data,
 139                                                                                                   unsigned int L, unsigned int H, unsigned int p){
 140   // coordonnes du point dans l'image
 141   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
 142   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;
 143   unsigned int pos = i*L +j ;
 144
 145   d_estim[ pos ] = ( d_data[ pos ] + p*d_lniv[ pos ])/(1+p) ;
 146
 147 }
 148
 149 __global__ void kernel_levelines_global_mem(unsigned int * img_in, unsigned int * img_out, unsigned int L, unsigned int H)
 150 {
 151   // coordonnes du point dans l'image
 152   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
 153   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;
 154   //unsigned int spos = threadIdx.x * blockDim.y + threadIdx.y ;
 155
 156   // nb de points par chemin
 157   int lpath =  PSIZE_J ;
 158   unsigned int ic, jc, zc, pos = i*L+j;
 159   int idpath, idpix ;
 160   unsigned int mse_min, mse_cur, val ;
 161   uint2 mse ;
 162
 163
 164   if((i>lpath)&&(i<H-lpath)&&(j>lpath)&&(j<L-lpath)){
 165         for( idpath=0; idpath < PSIZE_I ; idpath++) {
 166           ic = i ;
 167           jc = j ;
 168           pos = ic*L + jc ;
 169           zc = img_in[ pos ] ;
 170           mse.x = zc ;
 171           mse.y = zc*zc ;
 172           for( idpix=0; idpix < lpath-1 ; idpix++ ) {
 173                 ic += pathDi[idpath][idpix] ;
 174                 jc += pathDj[idpath][idpix] ;
 175                 pos = ic*L + jc ;
 176                 zc = img_in[ pos ] ;
 177                 mse.x += zc ;
 178                 mse.y += zc*zc ;
 179           }
 180           // critere de selection du chemin ( SUM_(X2) - SUM_(X)2 / lpath )
 181           // a ameliorer pour vitesse
 182           mse_cur = ( mse.y - ( mse.x / lpath ) * mse.x ) ;
 183           if (idpath == 0) {
 184                 mse_min = mse_cur ;
 185                 val = mse.x ;
 186           } else {
 187                 if ( mse_cur < mse_min )  {
 188                   mse_min = mse_cur ;
 189                   val = mse.x ;
 190                 }
 191           }
 192         }
 193         img_out[ i*L + j ] = val / lpath ;
 194   }
 195 }
 196
 197 __global__ void kernel_levelines_texture(unsigned int * img_out, unsigned int L, unsigned int H)
 198 {
 199   // coordonnes du point dans l'image
 200   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
 201   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;
 202   //unsigned int spos = threadIdx.x * blockDim.y + threadIdx.y ;
 203
 204   // nb de points par chemin
 205   int lpath =  PSIZE_J ;
 206   unsigned int ic, jc, zc ;
 207   int idpath, idpix ;
 208   unsigned int mse_min, mse_cur, val ;
 209   uint2 mse ;
 210
 211
 212   if((i>lpath)&&(i<H-lpath)&&(j>lpath)&&(j<L-lpath)){
 213         for( idpath=0; idpath < PSIZE_I ; idpath++) {
 214           ic = i ;
 215           jc = j ;
 216           zc = tex2D(tex_img_estim, j, i) ;
 217           mse.x = zc ;
 218           mse.y = zc*zc ;
 219           for( idpix=0; idpix < lpath-1 ; idpix++ ) {
 220                 ic += pathDi[idpath][idpix] ;
 221                 jc += pathDj[idpath][idpix] ;
 222                 zc = tex2D(tex_img_estim, jc, ic) ;
 223                 mse.x += zc ;
 224                 mse.y += zc*zc ;
 225           }
 226           // critere de selection du chemin ( SUM_(X2) - SUM_(X)2 / lpath )
 227           // a ameliorer pour vitesse
 228           mse_cur = ( mse.y - ( mse.x / lpath ) * mse.x ) ;
 229           if (idpath == 0) {
 230                 mse_min = mse_cur ;
 231                 val = mse.x ;
 232           } else {
 233                 if ( mse_cur < mse_min )  {
 234                   mse_min = mse_cur ;
 235                   val = mse.x ;
 236                 }
 237           }
 238         }
 239         img_out[ i*L + j ] = val / lpath ;
 240   }
 241 }
 242
 243
 244
 245
 246 __global__ void kernel_levelines_only_texture(unsigned int * img_out, unsigned int L, unsigned int H)
 247 {
 248   // coordonnes du point dans l'image
 249   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
 250   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;;
 251   //unsigned int spos = threadIdx.x * blockDim.y + threadIdx.y ;
 252
 253   // nb de points par chemin
 254   int lpath =  PSIZE_J ;
 255   unsigned int ic, jc ;
 256   int idpath, idpix ;
 257   unsigned int mse_min, mse_cur ;
 258
 259   //extern __shared__ uint2 mse[] ;
 260   uint2 mse ;
 261
 262   if((i>lpath)&&(i<H-lpath)&&(j>lpath)&&(j<L-lpath)){
 263         for( idpath=0; idpath < PSIZE_I ; idpath++) {
 264           ic = i ;
 265           jc = j ;
 266           mse.x = tex2D(tex_img_in, i, j) ;
 267           mse.y = tex2D(tex_img_in, i, j)*tex2D(tex_img_in, i, j) ;
 268           for( idpix=0; idpix < lpath-1 ; idpix++ ) {
 269                 ic += pathDi[idpath][idpix] ;
 270                 jc += pathDj[idpath][idpix] ;
 271                 mse.x += tex2D( tex_img_in, ic, jc ) ;
 272                 mse.y += tex2D( tex_img_in, ic, jc ) * tex2D( tex_img_in, ic, jc ) ;
 273           }
 274           // critere de selection du chemin ( SUM_(X2) - SUM_(X)2 / lpath )
 275           // a ameliorer pour vitesse
 276           mse_cur = ( mse.y - ( mse.x / lpath ) * mse.x ) ;
 277           if (idpath > 0) {
 278                 if ( mse_cur < mse_min )  {
 279                   mse_min = mse_cur ;
 280                 }
 281           } else {
 282                 mse_min = mse_cur ;
 283           }
 284         }
 285         img_out[ i*L + j ] = mse_min / lpath ;
 286   }
 287 }
 288
 289
 290 __global__ void kernel_trace_levelines(unsigned int * img_in, unsigned int * dir, unsigned int * img_out,
 291                                                                            unsigned int L, unsigned int H, unsigned int pas, unsigned int ng){
 292   // coordonnes du point dans l'image
 293   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
 294   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;;
 295
 296   // nb de points par chemin
 297   int lpath =  PSIZE_J ;
 298   unsigned int ic, jc, idpix ;
 299   unsigned int idpath  ;
 300
 301   img_out[ i*L + j ] = img_in[ i*L + j ] ;
 302
 303   if ( !(i%pas+j%pas)&&(i>lpath)&&(i<H-lpath)&&(j>lpath)&&(j<L-lpath) ){
 304         ic = i ;
 305         jc = j ;
 306         idpath = dir[ic*L+jc] ;
 307         img_out[ ic*L+jc ] = ng ;
 308         for ( idpix=0 ; idpix < lpath-1 ; idpix++ ){
 309           ic += pathDi[idpath][idpix] ;
 310           jc += pathDj[idpath][idpix] ;
 311           img_out[ ic*L + jc ] = ng ;
 312           }
 313   }
 314
 315 }