]> AND Private Git Repository - book_gpu.git/blob - BookGPU/Chapters/chapter4/code/convoGene8r.cu~
Logo AND Algorithmique Numérique Distribuée

Private GIT Repository
modif gillou
[book_gpu.git] / BookGPU / Chapters / chapter4 / code / convoGene8r.cu~
1 __global__ void kernel_convoGene8r( unsigned char  *output, int j_dim, int r)
2 {
3   int ic, jc ;
4   int L=2*r+1 ;
5   float outval0=0.0 ;
6   
7   // absolute coordinates of base point
8   int j = __umul24( blockIdx.x, blockDim.x ) + threadIdx.x ; 
9   int i = __umul24( blockIdx.y, blockDim.y) + threadIdx.y ; 
10
11   // convolution computation
12   for (ic=-r ; ic<=r ; ic++)
13         for(jc=-r ; jc<=r ; jc++)
14           outval0 += masque[ __umul24(ic,L)+jc+r ]
15                         *tex2D(tex_img_inc, j+jc, i+ic) ;
16
17   output[ __umul24(i, j_dim) + j ] = outval0 ;
18 }