climage/
gray.rs

1use super::context::ClContext;
2use super::imageproc::ImageProc;
3use ocl::{Kernel, Program};
4
5static KERNEL_SRC: &'static str = r#"
6__constant sampler_t sampler_const = CLK_NORMALIZED_COORDS_FALSE |
7                                    CLK_ADDRESS_NONE |
8                                    CLK_FILTER_NEAREST;
9__kernel void grayscale(
10    __read_only image2d_t src_image,
11    __write_only image2d_t dst_image) {
12    int2 gid = (int2)(get_global_id(0), get_global_id(1));
13    int2 size = get_image_dim(src_image);
14
15    if(all(gid < size)){
16        uint4 pixel = read_imageui(src_image, sampler_const, gid);
17        float4 color = convert_float4(pixel) / 255;
18        color.xyz = 0.2126 * color.x + 0.7152 * color.y + 0.0722 * color.z;
19        pixel = convert_uint4_rte(color * 255);
20        write_imageui(dst_image, gid, pixel);
21    }
22}
23"#;
24
25pub struct Gray<'a> {
26    context: &'a ClContext,
27    program: Program,
28    kernel: Option<Kernel>,
29}
30
31impl<'a> Gray<'a> {
32    pub fn new(context: &'a ClContext) -> Self {
33        let program = Program::builder()
34            .src(KERNEL_SRC)
35            .build(&context.context)
36            .unwrap();
37        Gray {
38            context,
39            program,
40            kernel: None,
41        }
42    }
43}
44
45impl<'a> ImageProc for Gray<'a> {
46    fn build_kernel(
47        &mut self,
48        src_img: &crate::ClImageBuffer,
49        dst_img: &mut crate::ClImageBuffer,
50    ) -> &mut Self {
51        let kernel = ocl::Kernel::builder()
52            .program(&self.program)
53            .name("grayscale")
54            .queue(self.context.queue.clone())
55            .global_work_size(src_img.dimensions())
56            .arg(&src_img.data)
57            .arg(&dst_img.data)
58            .build()
59            .unwrap();
60        self.kernel = Some(kernel);
61        self
62    }
63    fn run(&self) {
64        unsafe { self.kernel.as_ref().unwrap().enq().unwrap() };
65    }
66}