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
49
50
51
52
|
const constant float filter_flag = 0.111111f;
__kernel void
bench_copy_image(__read_only image2d_t src, __write_only image2d_t dst)
{
uint4 color = 0;
int2 coord;
int x = (int)get_global_id(0);
int y = (int)get_global_id(1);
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_CLAMP| CLK_FILTER_NEAREST;
coord.x = x;
coord.y = y;
color=read_imageui(src, sampler, coord);
write_imageui(dst, coord, color);
}
__kernel void
bench_filter_image(__read_only image2d_t src, __write_only image2d_t dst)
{
float4 color = 0;
int2 coord_00, coord_01, coord_02, coord_10, coord_11, coord_12, coord_20, coord_21, coord_22;
int x = (int)get_global_id(0);
int y = (int)get_global_id(1);
int x_sz = (int)get_global_size(0);
int y_sz = (int)get_global_size(1);
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_CLAMP| CLK_FILTER_NEAREST;
int x0 = x - 1; int x1 = x + 1;
int y0 = y - 1; int y1 = y + 1 ;
int x_left = (x > 0)?x0:x; int x_right = (x > x_sz - 2)?x:x1;
int y_top = (y > 0)?y0:y; int y_bottom = (y > y_sz - 2)?y:y1;
coord_00.x = x_left; coord_00.y = y_top;
coord_01.x = x; coord_01.y = y_top;
coord_02.x = x_right; coord_02.y = y_top;
coord_10.x = x_left; coord_10.y = y;
coord_11.x = x; coord_11.y = y;
coord_12.x = x_right; coord_12.y = y;
coord_20.x = x_left; coord_20.y = y_bottom;
coord_21.x = x; coord_21.y = y_bottom;
coord_22.x = x_right; coord_22.y = y_bottom;
color = convert_float4(read_imageui(src, sampler, coord_00)) + convert_float4(read_imageui(src, sampler, coord_01)) + convert_float4(read_imageui(src, sampler, coord_02))
+ convert_float4(read_imageui(src, sampler, coord_10)) + convert_float4(read_imageui(src, sampler, coord_11)) + convert_float4(read_imageui(src, sampler, coord_12))
+ convert_float4(read_imageui(src, sampler, coord_20)) + convert_float4(read_imageui(src, sampler, coord_21)) + convert_float4(read_imageui(src, sampler, coord_22));
write_imageui(dst, coord_11, convert_uint4(color * filter_flag));
}
|
