Game Development Reference
In-Depth Information
[ 2 2 j
]
[ 2 i + 1
2 j ]
[ i + j ]
[ j ]
[ i + 0.5
]
j + 0.5
2
,
[ 2 i + 1
2 j + 1 ]
[ 2 j + 1
]
2 i
2
(a)
(b)
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 ); if (tmp.w > 0) { sum += tmp; }
tmp = st(2*ix+1, 2*iy ); if (tmp.w > 0) { sum += tmp; }
tmp = st(2*ix, 2*iy+1); if (tmp.w > 0) { sum += tmp; }
tmp = st(2*ix+1, 2*iy+1); if (tmp.w > 0) { sum += tmp; }
if (sum.w > 0) sum /= sum.w;
dst(ix, iy) = sum;
}
Listing 5.3. Implementation of the restiction operation.
__global__ void interpolate( const gpu_plm2<float4> st_fine,
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 st = st_fine(ix, iy);
if (st.w < 1) {
st = make_float4 ( make_float3 (
tex2D (texST, 0.5f * (ix + 0.5f),
0.5f * (iy + 0.5f) )), 0);
}
dst(ix, iy) = st;
}
Listing 5.4. Implementation of the interpolation operation.