ch15

[book_gpu.git] / BookGPU / Chapters / chapter4 / code / convoGeneSh1.cu

diff --git a/BookGPU/Chapters/chapter4/code/convoGeneSh1.cu b/BookGPU/Chapters/chapter4/code/convoGeneSh1.cu
index 53d9ba45d73f44d6bc4cac3b1075098526afa29a..95eed34620180cc9b22dbd8f62ad6cb1c0eab390 100644 (file)
--- a/BookGPU/Chapters/chapter4/code/convoGeneSh1.cu
+++ b/BookGPU/Chapters/chapter4/code/convoGeneSh1.cu
@@ -11,7 +11,7 @@ __global__ void kernel_convoNonSepSh_8p(unsigned char *output, int j_dim, int r)
   int i = __umul24( blockIdx.y, blockDim.y) + threadIdx.y ;
   int j0= __umul24(blockIdx.x,blockDim.x)<<3 ;   // block's base point
   int idx = __umul24(i,j_dim) + j ;              // absolute index
   int i = __umul24( blockIdx.y, blockDim.y) + threadIdx.y ;
   int j0= __umul24(blockIdx.x,blockDim.x)<<3 ;   // block's base point
   int idx = __umul24(i,j_dim) + j ;              // absolute index

-  int idrow = threadIdx.y*(bdimX+k-1) ;          // line's offset in sh mem
+  int idrow = threadIdx.y*(bdimX+k-1) ;    // line's offset in sh mem

   
   // shared memory declaration
   extern __shared__ unsigned char roi8p[];
   
   // shared memory declaration
   extern __shared__ unsigned char roi8p[];


@@ -46,7 +46,7 @@ __global__ void kernel_convoNonSepSh_8p(unsigned char *output, int j_dim, int r)
        for( jc=0 ; jc<k ; jc++)
          {
                int baseRoi = __umul24(ic+threadIdx.y,(bdimX+k-1)) + jc+tidX ;
        for( jc=0 ; jc<k ; jc++)
          {
                int baseRoi = __umul24(ic+threadIdx.y,(bdimX+k-1)) + jc+tidX ;

-               float valMask = masque[ __umul24(ic,k)+jc ] ;
+               float valMask = mask[ __umul24(ic,k)+jc ] ;

                outval0 += valMask*roi8p[ baseRoi ] ;
                outval1 += valMask*roi8p[ baseRoi +1 ] ;
                outval2 += valMask*roi8p[ baseRoi +2 ] ;
                outval0 += valMask*roi8p[ baseRoi ] ;
                outval1 += valMask*roi8p[ baseRoi +1 ] ;
                outval2 += valMask*roi8p[ baseRoi +2 ] ;


@@ -65,5 +65,5 @@ __global__ void kernel_convoNonSepSh_8p(unsigned char *output, int j_dim, int r)
   output[ idx++ ] = outval4 ;
   output[ idx++ ] = outval5 ;
   output[ idx++ ] = outval6 ;
   output[ idx++ ] = outval4 ;
   output[ idx++ ] = outval5 ;
   output[ idx++ ] = outval6 ;

-  output[ idx++ ] = outval7 ;
+  output[ idx   ] = outval7 ;

 }
 }