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}