test tex

[snake_gpu.git] / src / lib_kernels_cumuls.cu

diff --git a/src/lib_kernels_cumuls.cu b/src/lib_kernels_cumuls.cu
index 641a0ff004a02a5c9bf10c308c0318962680c3c6..070a85f1271eaad83db644a5b59137674b152c45 100644 (file)
--- a/src/lib_kernels_cumuls.cu
+++ b/src/lib_kernels_cumuls.cu
@@ -1,5 +1,5 @@
 
 

-
+texture<int, 2, cudaReadModeElementType> tex_img_in ;

 
 __global__ void calcul_cumuls_gpu(unsigned short * img, t_cumul_x * cumul_x, t_cumul_x2 * cumul_x2, uint32 h, uint32 l, uint64 * gsomblocs, unsigned int N_BLOCS_LIGNE, unsigned int baseline, unsigned int nb_lines )
 {
 
 __global__ void calcul_cumuls_gpu(unsigned short * img, t_cumul_x * cumul_x, t_cumul_x2 * cumul_x2, uint32 h, uint32 l, uint64 * gsomblocs, unsigned int N_BLOCS_LIGNE, unsigned int baseline, unsigned int nb_lines )
 {


@@ -13,13 +13,15 @@ __global__ void calcul_cumuls_gpu(unsigned short * img, t_cumul_x * cumul_x, t_c
   int pos = num_ligne*l + til ;                     //indice du pixel dans la 'tranche' courante
   int pos_img = pos + baseline*l ;                  //indice correspondant dans l'image originale
   int id, stride=1;
   int pos = num_ligne*l + til ;                     //indice du pixel dans la 'tranche' courante
   int pos_img = pos + baseline*l ;                  //indice correspondant dans l'image originale
   int id, stride=1;

+  unsigned int z ;

   
   extern __shared__ tcumuls sdata[];
   //chargement en smem avec complement de 0  
   id = CFI(tib);
   if (til < l){
   
   extern __shared__ tcumuls sdata[];
   //chargement en smem avec complement de 0  
   id = CFI(tib);
   if (til < l){

-           sdata[id].x = img[ pos_img ] ;
-           sdata[id].x2 = img[ pos_img ]*img[ pos_img ] ;
+       z = tex2D(tex_img_in, til, num_ligne) ;
+           sdata[id].x = z ;
+           sdata[id].x2 = z*z ;

         } else {
           sdata[id].x = 0 ;
           sdata[id].x2 = 0 ;
         } else {
           sdata[id].x = 0 ;
           sdata[id].x2 = 0 ;