X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/lniv_gpu.git/blobdiff_plain/1a59c41c272e22861e759e660a64578b47147d85..cb7f1415a3a53d241784928e2a5d291f79a8e984:/levelines_kernels.cu

diff --git a/levelines_kernels.cu b/levelines_kernels.cu
index 491c600..754f3bf 100644
--- a/levelines_kernels.cu
+++ b/levelines_kernels.cu
@@ -6,65 +6,65 @@
 __constant__ int pathDi[PSIZE_I][PSIZE_J-1] =
   {
     // Q1
-	{  -1, -1, -1},       // 
-	{  -1, -1, -1},       //  
-	{  -1, -1, -1},       // 
-	{  -1, -1, -1},       // 
-	{  -1,  0,  1},       // 
-	{   0, -1,  0},
+	{  -1, -1, -1},       // 90
+	{  -1, -1, -1},       // 75
+	{  -1, -1, -1},       // 60
+	{  -1, -1, -1},       // 45
+	{  -1,  0, -1},       // 30
+	{   0, -1,  0},       // 15
 	// Q4
-	{   0,  0,  0},       // 
-	{   0,  1,  1},       //  
-	{   1,  0,  1},       // 
-	{   1,  1,  1},       // 
-	{   1,  1,  1},       // 
-	{   1,  1,  1},
+	{   0,  0,  0},       // 0
+	{   0,  1,  0},       // 345 
+	{   1,  0,  1},       // 330
+	{   1,  1,  1},       // 315
+	{   1,  1,  1},       // 300
+	{   1,  1,  1},       // 285
 	// Q3
-	{   1,  1,  1},       // 
-	{   1,  1,  1},       //  
-	{   1,  1,  1},       // 
-	{   1,  1,  1},       // 
-	{   1,  0, -1},       // 
-	{   0,  1,  0},
+	{   1,  1,  1},       // 270
+	{   1,  1,  1},       // 255 
+	{   1,  1,  1},       // 240
+	{   1,  1,  1},       // 225
+	{   1,  0,  1},       // 210
+	{   0,  1,  0},       // 195
 	// Q2
-	{   0,  0,  0},       // 
-	{   0, -1,  0},       //  
-	{  -1,  0, -1},       // 
-	{  -1, -1, -1},       // 
-	{  -1, -1, -1},       // 
-	{  -1, -1, -1}
+	{   0,  0,  0},       // 180
+	{   0, -1,  0},       // 165 
+	{  -1,  0, -1},       // 150
+	{  -1, -1, -1},       // 135
+	{  -1, -1, -1},       // 120
+	{  -1, -1, -1}        // 105
   } ;     // 
   
 __constant__ int pathDj[PSIZE_I][PSIZE_J-1] =
   {
 	// Q1
-	{  0,  0,  0},       // 
-	{  0,  1,  0},
-	{  1,  0,  1},  
-	{  1,  1,  1},  
-	{  1,  1,  1},  
-	{  1,  1,  1},
+	{  0,  0,  0},       // 90
+	{  0,  1,  0},       // 75
+	{  1,  0,  1},       // 60
+	{  1,  1,  1},       // 45
+	{  1,  1,  1},       // 30 
+	{  1,  1,  1},       // 15
 	// Q4
-	{  1,  1,  1},       // 
-	{  1,  1,  1},
-	{  1,  1,  1},  
-	{  1,  1,  1},  
-	{  1,  0, -1},  
-	{  0,  1,  0},
-	// Q3
-	{  0,  0,  0},       // 
-	{  0, -1,  0},
-	{ -1,  0, -1},  
-	{ -1, -1, -1},  
-	{ -1, -1, -1},  
-	{ -1, -1, -1},
+	{  1,  1,  1},       // 0
+	{  1,  1,  1},       // 345
+	{  1,  1,  1},       // 330
+	{  1,  1,  1},       // 315  
+	{  1,  0,  1},       // 300
+	{  0,  1,  0},       // 285
+	// Q3 
+	{  0,  0,  0},       // 270
+	{  0, -1,  0},       // 255
+	{ -1,  0, -1},       // 240
+	{ -1, -1, -1},       // 225
+	{ -1, -1, -1},       // 210
+	{ -1, -1, -1},       // 195
 	// Q2
-	{ -1, -1, -1},       // 
-	{ -1, -1, -1},
-	{ -1, -1, -1},  
-	{ -1, -1, -1},  
-	{ -1,  0,  1},  
-	{  0, -1,  0}
+	{ -1, -1, -1},       // 180
+	{ -1, -1, -1},       // 165 
+	{ -1, -1, -1},       // 150
+	{ -1, -1, -1},       // 135
+	{ -1,  0, -1},       // 120 
+	{  0, -1,  0}        // 105
   } ;     
 
   
@@ -74,6 +74,72 @@ texture<int, 2, cudaReadModeElementType> tex_img_estim ;
 texture<int, 2, cudaReadModeElementType> tex_img_lniv ;
 
 
+__constant__ float tangente[] = {0.000, 0.268, 0.577, 1.000} ;
+
+__global__ void kernel_calcul_paths( int2 * d_paths, unsigned int r){
+
+  unsigned int idpath = 0 ;
+  int ic, jc, iprec, jprec ;
+  float offset = 0.5 ;
+  unsigned int basepath = 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 ;
+	  }
+	  iprec = ic ;
+	}
+	idpath++ ;
+  }
+  // Q1 sup
+  for (int a=2 ; a>0 ; a--){         // les 2 angles 60 et 75 
+	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 ;
+	  }
+	  jprec = jc ;
+	}
+	idpath++ ;
+  }
+  
+  // Q2
+  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 ;
+	  }
+	idpath++ ;
+  }
+
+  // Q3
+  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 ;
+	  }
+	idpath++ ;
+  }
+
+  // Q4
+  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 ;
+	  }
+	idpath++ ;
+  }
+}
+
+
 __global__ void kernel_init_estim_from_img_in(unsigned int * d_estim, unsigned int L, unsigned int H, unsigned int r){
   // coordonnes du point dans l'image
   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
@@ -96,6 +162,17 @@ __global__ void kernel_init_estim_from_img_in(unsigned int * d_estim, unsigned i
   */
 }
 
+__global__ void kernel_neutre_estim_from_img_in( unsigned int * d_estim, unsigned int L, unsigned int H ){
+  // coordonnes 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 pos = i*L +j ;
+  
+  d_estim[ pos ] = tex2D(tex_img_in, j, i) ;
+ 
+
+}
+
 __global__ void kernel_init_estim_from_img_in_global_mem(unsigned int * d_data, unsigned int * d_estim,
 														 unsigned int L, unsigned int H, unsigned int r){
   // coordonnes du point dans l'image
@@ -209,7 +286,7 @@ __global__ void kernel_levelines_texture(unsigned int * img_out, unsigned int L,
   uint2 mse ;
   
   
-  if((i>lpath)&&(i<H-lpath)&&(j>lpath)&&(j<L-lpath)){
+  if( (i >=lpath )&&( i<= (H-lpath) )&&( j>=lpath )&&( j<=(L-lpath) ) ){
 	for( idpath=0; idpath < PSIZE_I ; idpath++) {
 	  ic = i ;
 	  jc = j ;
@@ -225,7 +302,7 @@ __global__ void kernel_levelines_texture(unsigned int * img_out, unsigned int L,
 	  }
 	  // critere de selection du chemin ( SUM_(X2) - SUM_(X)2 / lpath )
 	  // a ameliorer pour vitesse
-	  mse_cur = ( mse.y - ( mse.x / lpath ) * mse.x ) ;
+	  mse_cur = ( mse.y - ( mse.x*mse.x / lpath ) ) ;
 	  if (idpath == 0) {
 		mse_min = mse_cur ;
 		val = mse.x ;
@@ -238,56 +315,61 @@ __global__ void kernel_levelines_texture(unsigned int * img_out, unsigned int L,
 	}
 	img_out[ i*L + j ] = val / lpath ; 
   }
+  else img_out[ i*L + j ] = 0 ;
 }
 
 
-
-
-__global__ void kernel_levelines_only_texture(unsigned int * img_out, unsigned int L, unsigned int H)
+__global__ void kernel_dir_levelines_texture(unsigned int * dir, unsigned int L, unsigned int H)
 {
   // coordonnes 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 j = blockIdx.y*blockDim.y + threadIdx.y;
   //unsigned int spos = threadIdx.x * blockDim.y + threadIdx.y ;
   
   // nb de points par chemin
   int lpath =  PSIZE_J ;
-  unsigned int ic, jc ;
+  unsigned int ic, jc, zc ;
   int idpath, idpix ;
-  unsigned int mse_min, mse_cur ;
-  
-  //extern __shared__ uint2 mse[] ;
+  unsigned int mse_min, mse_cur, val ;
   uint2 mse ;
   
-  if((i>lpath)&&(i<H-lpath)&&(j>lpath)&&(j<L-lpath)){
+  
+  if( (i >=lpath )&&( i<= (H-lpath) )&&( j>=lpath )&&( j<=(L-lpath) ) ){
 	for( idpath=0; idpath < PSIZE_I ; idpath++) {
 	  ic = i ;
 	  jc = j ;
-	  mse.x = tex2D(tex_img_in, i, j) ;
-	  mse.y = tex2D(tex_img_in, i, j)*tex2D(tex_img_in, i, j) ;
+	  zc = tex2D(tex_img_estim, j, i) ;
+	  mse.x = zc ;
+	  mse.y = zc*zc ;
 	  for( idpix=0; idpix < lpath-1 ; idpix++ ) {
 		ic += pathDi[idpath][idpix] ;
 		jc += pathDj[idpath][idpix] ;
-		mse.x += tex2D( tex_img_in, ic, jc ) ;
-		mse.y += tex2D( tex_img_in, ic, jc ) * tex2D( tex_img_in, ic, jc ) ; 
+		zc = tex2D(tex_img_estim, jc, ic) ;
+		mse.x += zc ;
+		mse.y += zc*zc ; 
 	  }
 	  // critere de selection du chemin ( SUM_(X2) - SUM_(X)2 / lpath )
 	  // a ameliorer pour vitesse
-	  mse_cur = ( mse.y - ( mse.x / lpath ) * mse.x ) ;
-	  if (idpath > 0) {
+	  mse_cur = ( mse.y - ( mse.x*mse.x / lpath ) ) ;
+	  if (idpath == 0) {
+		mse_min = mse_cur ;
+		val = idpath ;
+	  } else {
 		if ( mse_cur < mse_min )  {
 		  mse_min = mse_cur ;
+		  val = idpath ; 
 		}
-	  } else {
-		mse_min = mse_cur ;
-	  }
+	  } 
 	}
-	img_out[ i*L + j ] = mse_min / lpath ; 
+	dir[ i*L + j ] = val ; 
   }
+  
 }
 
 
-__global__ void kernel_trace_levelines(unsigned int * img_in, unsigned int * dir, unsigned int * img_out,
+
+
+__global__ void kernel_trace_levelines_texture(unsigned int * dir, unsigned int * img_out,
 									   unsigned int L, unsigned int H, unsigned int pas, unsigned int ng){
   // coordonnes du point dans l'image
   unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
@@ -298,9 +380,9 @@ __global__ void kernel_trace_levelines(unsigned int * img_in, unsigned int * dir
   unsigned int ic, jc, idpix ;
   unsigned int idpath  ;
 
-  img_out[ i*L + j ] = img_in[ i*L + j ] ;
+  img_out[ i*L + j ] = tex2D(tex_img_estim, j, i ) ;
   
-  if ( !(i%pas+j%pas)&&(i>lpath)&&(i<H-lpath)&&(j>lpath)&&(j<L-lpath) ){
+  if ( !(i%pas+j%pas)&&(i>=lpath)&&(i<=H-lpath)&&(j>=lpath)&&(j<=L-lpath) ){
 	ic = i ;
 	jc = j ;
 	idpath = dir[ic*L+jc] ;