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 ;
* 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 ;
}
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 ;
}
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++ ;
}
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++ ;
}
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++ ;
}
+
}
/**
* 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;
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) ;
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 ;
}
-/**
- *
- * \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é
* 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 ){
}
}
+*/
