X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/lniv_gpu.git/blobdiff_plain/7cc36a5689750ee7d77b959362edc65948a79427..HEAD:/levelines_kernels.cu?ds=inline diff --git a/levelines_kernels.cu b/levelines_kernels.cu index 6838e83..6d616c6 100644 --- a/levelines_kernels.cu +++ b/levelines_kernels.cu @@ -78,7 +78,7 @@ __constant__ float tangente[] = {0.000, 0.268, 0.577, 1.000} ; texture tex_img_in ; texture tex_img_estim ; texture tex_img_lniv ; -texture tex_paths ; +texture tex_paths ; @@ -95,20 +95,24 @@ texture tex_paths ; * 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 ; + 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 ){ - 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 ; } @@ -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 ){ - 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 ; } @@ -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> 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++ ; } @@ -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> 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++ ; } @@ -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> 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++ ; } + } /** @@ -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". - * 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) { - // 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; @@ -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 ; + 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) ; @@ -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++ ) { - 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 ; @@ -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é @@ -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 */ +/* __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 } } +*/