GPU Pro 4

5. Coherence-Enhancing Filtering on the GPU 237

[

i

j

]

[

2j + 1

2i

]

[

2i + 1

[i + j]

2j

]

[

2i + 1

2j + 1

]

[

i + 0.5

2

,

]

[

2i

2j

(a) (b)

]

j + 0.5

2

Figure 5.7. Illustration of the restriction and interpolation operations. (a) Restriction

collapses four pixels into a single pixel by averaging those that are boundary pixels.

(b) Interpolation is performed using bilinear interpolation.

__global__ void restrict ( const gpu_plm2 <float4 > st,

gpu_plm2 < float4 > dst)

{

const int ix = blockDim .x * blockIdx .x + threadIdx.x;

const int iy = blockDim .y * blockIdx .y + threadIdx.y;

if (ix >= dst.w || iy >= dst.h) return;

float4 sum = make_float4(0);

float4 tmp;

tmp = st (2*ix , 2*iy ...