version opérationnelle

[lniv_gpu.git] / levelines_kernels.cu

diff --git a/levelines_kernels.cu b/levelines_kernels.cu
index 6838e8322a7e2e786910c26cbaa0a5132f5ccd09..6d616c66c3c1f11e7810da00329c8d9c2095e5e4 100644 (file)
--- a/levelines_kernels.cu
+++ b/levelines_kernels.cu
@@ -78,7 +78,7 @@ __constant__ float tangente[] = {0.000, 0.268, 0.577, 1.000} ;
 texture<int, 2, cudaReadModeElementType> tex_img_in ;
 texture<int, 2, cudaReadModeElementType> tex_img_estim ;
 texture<int, 2, cudaReadModeElementType> tex_img_lniv ;
 texture<int, 2, cudaReadModeElementType> tex_img_in ;
 texture<int, 2, cudaReadModeElementType> tex_img_estim ;
 texture<int, 2, cudaReadModeElementType> tex_img_lniv ;

-texture<int2, 2, cudaReadModeElementType> tex_paths ;
+texture<ushort, 2, cudaReadModeElementType> tex_paths ;

 
 
 
 
 
 


@@ -95,20 +95,24 @@ texture<int2, 2, cudaReadModeElementType> tex_paths ;
  * considérés pour le calcul de chemins (float tangente[]).
  * 
  */
  * considérés pour le calcul de chemins (float tangente[]).
  * 
  */

-__global__ void kernel_calcul_paths( int2 * d_paths, unsigned int r){
+__global__ void kernel_calcul_paths( ushort * d_paths, unsigned int r){

 
   unsigned int idpath = 0 ;
   int ic, jc, iprec, jprec ;
   float offset = 0.5 ;
   unsigned int basepath = 0 ;
 
   unsigned int idpath = 0 ;
   int ic, jc, iprec, jprec ;
   float offset = 0.5 ;
   unsigned int basepath = 0 ;

+  char MSQ, LSQ ;

 
   // Q1 inf
   for (int a=0 ; a< 4 ; a++){        // les 4 angles 0,15,30 et 45
        for (int p=0 ; p< r ; p++){      // les r points
          ic = r-1 - floor(tangente[a]*p + offset) ;
          if ( p > 0 ){
 
   // Q1 inf
   for (int a=0 ; a< 4 ; a++){        // les 4 angles 0,15,30 et 45
        for (int p=0 ; p< r ; p++){      // les r points
          ic = r-1 - floor(tangente[a]*p + offset) ;
          if ( p > 0 ){

-               d_paths[idpath*(r-1)+p-1].x = ic - iprec ;
-               d_paths[idpath*(r-1)+p-1].y = 1 ;
+               MSQ = ic - iprec ;
+               LSQ = 1 ;
+               //d_paths[idpath*(r-1)+p-1].x = ic - iprec ;
+               //d_paths[idpath*(r-1)+p-1].y = 1 ;
+               d_paths[idpath*(r-1)+p-1] = ((short)MSQ << 8) | LSQ ;

          }
          iprec = ic ;
        }
          }
          iprec = ic ;
        }


@@ -119,8 +123,11 @@ __global__ void kernel_calcul_paths( int2 * d_paths, unsigned int r){
        for (int p=0 ; p< r ; p++){      // les r points
          jc = floor(tangente[a]*p + offset) ; 
          if ( p > 0 ){
        for (int p=0 ; p< r ; p++){      // les r points
          jc = floor(tangente[a]*p + offset) ; 
          if ( p > 0 ){

-               d_paths[idpath*(r-1)+p-1].x = -1 ;
-               d_paths[idpath*(r-1)+p-1].y = jc - jprec ;
+               MSQ = -1 ;
+               LSQ = jc - jprec ;
+               d_paths[idpath*(r-1)+p-1] = ((short)MSQ << 8) | LSQ ;
+               //d_paths[idpath*(r-1)+p-1].x = -1 ;
+               //d_paths[idpath*(r-1)+p-1].y = jc - jprec ;

          }
          jprec = jc ;
        }
          }
          jprec = jc ;
        }


@@ -131,8 +138,11 @@ __global__ void kernel_calcul_paths( int2 * d_paths, unsigned int r){
   basepath += 6 ;
   for (int a=0 ; a< 6 ; a++){         // les 6 angles 90,105,120,135,150,165
        for (int p=0 ; p<r-1 ; p++){      // les r points
   basepath += 6 ;
   for (int a=0 ; a< 6 ; a++){         // les 6 angles 90,105,120,135,150,165
        for (int p=0 ; p<r-1 ; p++){      // les r points

-         d_paths[idpath*(r-1)+p].x = -d_paths[(idpath - basepath)*(r-1)+p].y ;
-         d_paths[idpath*(r-1)+p].y =  d_paths[(idpath - basepath)*(r-1)+p].x ;
+         MSQ = - ( d_paths[(idpath - basepath)*(r-1)+p] & 0x00FF ) ;
+         LSQ = ( d_paths[(idpath - basepath)*(r-1)+p] >> 8 ) ;
+         d_paths[idpath*(r-1)+p-1] = ((short)MSQ << 8) | LSQ ;
+         //d_paths[idpath*(r-1)+p].x = -d_paths[(idpath - basepath)*(r-1)+p].y ;
+         //d_paths[idpath*(r-1)+p].y =  d_paths[(idpath - basepath)*(r-1)+p].x ;

          }
        idpath++ ;
   }
          }
        idpath++ ;
   }


@@ -141,8 +151,11 @@ __global__ void kernel_calcul_paths( int2 * d_paths, unsigned int r){
   basepath += 6 ;
   for (int a=0 ; a< 6 ; a++){         // les 6 angles 180,195,210,225,240,255
        for (int p=0 ; p<r-1 ; p++){      // les r points
   basepath += 6 ;
   for (int a=0 ; a< 6 ; a++){         // les 6 angles 180,195,210,225,240,255
        for (int p=0 ; p<r-1 ; p++){      // les r points

-         d_paths[idpath*(r-1)+p].x = -d_paths[(idpath - basepath)*(r-1)+p].x ;
-         d_paths[idpath*(r-1)+p].y = -d_paths[(idpath - basepath)*(r-1)+p].y ;
+         MSQ = - ( d_paths[(idpath - basepath)*(r-1)+p] >> 8 ) ;
+         LSQ = - ( d_paths[(idpath - basepath)*(r-1)+p] & 0x00FF ) ;
+         d_paths[idpath*(r-1)+p-1] = ((short)MSQ << 8) | LSQ ;
+         //d_paths[idpath*(r-1)+p].x = -d_paths[(idpath - basepath)*(r-1)+p].x ;
+         //d_paths[idpath*(r-1)+p].y = -d_paths[(idpath - basepath)*(r-1)+p].y ;

          }
        idpath++ ;
   }
          }
        idpath++ ;
   }


@@ -151,11 +164,15 @@ __global__ void kernel_calcul_paths( int2 * d_paths, unsigned int r){
   basepath += 6 ;
   for (int a=0 ; a< 6 ; a++){         // les 6 angles 270,285,300,315,330,345
        for (int p=0 ; p<r-1 ; p++){      // les r points
   basepath += 6 ;
   for (int a=0 ; a< 6 ; a++){         // les 6 angles 270,285,300,315,330,345
        for (int p=0 ; p<r-1 ; p++){      // les r points

-         d_paths[idpath*(r-1)+p].x =  d_paths[(idpath - basepath)*(r-1)+p].y ;
-         d_paths[idpath*(r-1)+p].y = -d_paths[(idpath - basepath)*(r-1)+p].x ;
+         MSQ = d_paths[(idpath - basepath)*(r-1)+p] & 0x00FF ;
+         LSQ = - ( d_paths[(idpath - basepath)*(r-1)+p] >> 8 ) ;
+         d_paths[idpath*(r-1)+p-1] = ((short)MSQ << 8) | LSQ ;
+         //d_paths[idpath*(r-1)+p].x =  d_paths[(idpath - basepath)*(r-1)+p].y ;
+         //d_paths[idpath*(r-1)+p].y = -d_paths[(idpath - basepath)*(r-1)+p].x ;

          }
        idpath++ ;
   }
          }
        idpath++ ;
   }

+

 }
 
 /**
 }
 
 /**


@@ -368,12 +385,15 @@ __global__ void kernel_estim_next_step_global_mem(unsigned int * d_estim, unsign
  * Execution sur des blocs de threads 2D et une grille 2D
  * selon les dimensions de l'image.
  * L'image d'entrée doit être au préalable en mémoire texture pointée par "tex_img_estim".
  * Execution sur des blocs de threads 2D et une grille 2D
  * selon les dimensions de l'image.
  * L'image d'entrée doit être au préalable en mémoire texture pointée par "tex_img_estim".

- * Les matrices des chemins sont, elles, pointées par "tex_paths"
+ * Les matrices des chemins sont, elles, préalablement chargées en SHMEM depuis la texture"

  * Cette version ne fournit pas les indices des chemins pour les tracé éventuel des lniv.
  */
 __global__ void kernel_levelines_texture(unsigned int * img_out, unsigned int L, unsigned int H, unsigned int r)
 {
  * Cette version ne fournit pas les indices des chemins pour les tracé éventuel des lniv.
  */
 __global__ void kernel_levelines_texture(unsigned int * img_out, unsigned int L, unsigned int H, unsigned int r)
 {

-  // coordonnes du point dans l'image
+  // coordonnees du point dans le bloc
+  unsigned int iib = threadIdx.x ;
+  unsigned int jib = threadIdx.y ;
+  // coordonnees du point dans l'image

   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;
     
   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
   unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;
     


@@ -383,7 +403,18 @@ __global__ void kernel_levelines_texture(unsigned int * img_out, unsigned int L,
   int idpath, idpix ;
   unsigned int mse_min, mse_cur, val ;
   uint2 mse ;
   int idpath, idpix ;
   unsigned int mse_min, mse_cur, val ;
   uint2 mse ;

+  short texVal ;
+

   
   

+  extern __shared__ short shPath[] ;
+
+  unsigned int absPos = jib*8 + iib ;
+  if ( absPos < PSIZE_I ){
+       for ( idpix = 0; idpix < lpath-1; idpix++){
+         shPath[ idpix*24 + absPos ] = tex2D(tex_paths, idpix, absPos) ; 
+       }
+       syncthreads() ;
+  }

   
   if((i>=lpath)&&(i<=H-lpath)&&(j>=lpath)&&(j<=L-lpath)){
        z = tex2D(tex_img_estim, j, i) ;
   
   if((i>=lpath)&&(i<=H-lpath)&&(j>=lpath)&&(j<=L-lpath)){
        z = tex2D(tex_img_estim, j, i) ;


@@ -393,8 +424,9 @@ __global__ void kernel_levelines_texture(unsigned int * img_out, unsigned int L,
          mse.x = z ;
          mse.y = z*z ;
          for( idpix=0; idpix < lpath-1 ; idpix++ ) {
          mse.x = z ;
          mse.y = z*z ;
          for( idpix=0; idpix < lpath-1 ; idpix++ ) {

-               ic += tex2D(tex_paths, idpix, idpath).x ;
-               jc += tex2D(tex_paths, idpix, idpath).y ;
+               texVal = shPath[ idpix*24 + idpath ] ;
+               ic += (char)(texVal>>8) ;
+               jc += (char)(texVal) ;

                zc = tex2D(tex_img_estim, jc, ic) ;
                mse.x += zc ;
                mse.y += zc*zc ; 
                zc = tex2D(tex_img_estim, jc, ic) ;
                mse.x += zc ;
                mse.y += zc*zc ; 


@@ -412,73 +444,6 @@ __global__ void kernel_levelines_texture(unsigned int * img_out, unsigned int L,
 }
 
 
 }
 
 

-/**
- *
- * \brief determine les lniv en chaque point de l'image
- * \author zulu - AND
- *
- * \param[in] L         Largeur de l'image
- * \param[in] H         Hauteur de l'image
- * \param[in] r         longueur des segments
- *
- * \param[out] img_out  image des lniv 
- *
- * Execution sur des blocs de threads 2D et une grille 2D
- * selon les dimensions de l'image.
- * L'image d'entrée doit être au préalable en mémoire texture pointée par "tex_img_estim".
- * Les matrices des chemins sont, elles, pointées par "tex_paths"
- * Cette version ne fournit pas les indices des chemins pour les tracé éventuel des lniv.
- * Cette version tente d'utiliser la shared memory pour compenser la baisse de perf due aux chemins
- * paramétrables non constants.
- */
-
-__global__ void kernel_levelines_texture_smem(unsigned int * img_out, unsigned int L, unsigned int H, unsigned int r)
-{
-  // coordonnées du point dans le bloc
-  unsigned int iib = threadIdx.x ;
-  unsigned int jib = threadIdx.y ;
-  // coordonnes du point dans l'image
-  unsigned int i = blockIdx.x*blockDim.x + iib ;
-  unsigned int j = blockIdx.y*blockDim.y + jib ;
-    
-  // nb de points par chemin
-  int lpath =  r ;
-  int ic, jc ;
-  int idpath, idpix ;
-  unsigned int val, mse_cur, mse_min, z, zc ;
-  uint2 mse_data ;
-
-  //__shared__ unsigned int val_img[16*16] ;
-
-  //val_img[jib*16+iib] = tex2D(tex_img_estim, j, i) ;
-  
-  if((i>=lpath)&&(i<=H-lpath)&&(j>=lpath)&&(j<=L-lpath)){
-       z = tex2D(tex_img_estim, j, i) ;
-       for( idpath=0; idpath < PSIZE_I ; idpath++) {
-         ic = i ;
-         jc = j ;
-         mse_data.x = z ;
-         mse_data.y = z*z ;
-         mse_min = mse_data.y - mse_data.x/lpath*mse_data.y ;
-         for( idpix=0; idpix < lpath-1 ; idpix++ ) {
-               ic += tex2D(tex_paths, idpix, idpath).x ;
-               jc += tex2D(tex_paths, idpix, idpath).y ;
-               zc = tex2D(tex_img_estim, jc, ic) ;
-               mse_data.x += zc ;
-               mse_data.y += zc*zc ; 
-         }
-         // critere de selection du chemin ( SUM_(X2) - SUM_(X)2 / lpath )
-         // TODO cherchera  ameliorer pour vitesse
-         mse_cur = ( mse_data.y - ( mse_data.x / lpath ) * mse_data.x ) ;
-         if ( mse_cur < mse_min ){
-               mse_min = mse_cur ;
-               val = mse_data.x ;
-         } 
-       }
-       img_out[ i*L + j ] = val / lpath ; 
-  }
-}
-

 /**
  *
  * \brief trace les segments sur un maillage carré
 /**
  *
  * \brief trace les segments sur un maillage carré


@@ -498,6 +463,7 @@ __global__ void kernel_levelines_texture_smem(unsigned int * img_out, unsigned i
  * execution sur des blocs de threads 2D et une grille 2D
  * selon les dimensions de l'image
  */
  * execution sur des blocs de threads 2D et une grille 2D
  * selon les dimensions de l'image
  */

+/*

 __global__ void kernel_trace_levelines(unsigned int * img_in, unsigned int * dir, unsigned int * img_out,
                                                                           unsigned int L, unsigned int H, unsigned int pas, unsigned int ng,
                                                                           unsigned int r ){
 __global__ void kernel_trace_levelines(unsigned int * img_in, unsigned int * dir, unsigned int * img_out,
                                                                           unsigned int L, unsigned int H, unsigned int pas, unsigned int ng,
                                                                           unsigned int r ){


@@ -525,3 +491,4 @@ __global__ void kernel_trace_levelines(unsigned int * img_in, unsigned int * dir
   }
   
 }
   }
   
 }

+*/