weightp.cl
1.68 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
/* Weightp filter a downscaled image into a temporary output buffer.
* This kernel is launched once for each scale.
*
* Launch dimensions: width x height (in pixels)
*/
kernel void weightp_scaled_images( read_only image2d_t in_plane,
write_only image2d_t out_plane,
uint offset,
uint scale,
uint denom )
{
int gx = get_global_id( 0 );
int gy = get_global_id( 1 );
uint4 input_val;
uint4 output_val;
input_val = read_imageui( in_plane, sampler, (int2)(gx, gy));
output_val = (uint4)(offset) + ( ( ((uint4)(scale)) * input_val ) >> ((uint4)(denom)) );
write_imageui( out_plane, (int2)(gx, gy), output_val );
}
/* Weightp filter for the half-pel interpolated image
*
* Launch dimensions: width x height (in pixels)
*/
kernel void weightp_hpel( read_only image2d_t in_plane,
write_only image2d_t out_plane,
uint offset,
uint scale,
uint denom )
{
int gx = get_global_id( 0 );
int gy = get_global_id( 1 );
uint input_val;
uint output_val;
input_val = read_imageui( in_plane, sampler, (int2)(gx, gy)).s0;
//Unpack
uint4 temp;
temp.s0 = input_val & 0x00ff; temp.s1 = (input_val >> 8) & 0x00ff;
temp.s2 = (input_val >> 16) & 0x00ff; temp.s3 = (input_val >> 24) & 0x00ff;
temp = (uint4)(offset) + ( ( ((uint4)(scale)) * temp ) >> ((uint4)(denom)) );
//Pack
output_val = temp.s0 | (temp.s1 << 8) | (temp.s2 << 16) | (temp.s3 << 24);
write_imageui( out_plane, (int2)(gx, gy), output_val );
}