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
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108

use crate::{Error, Node, OpenCL};
use min_cl::api::{
    build_program, create_kernels_in_program, create_program_with_source, release_mem_object,
    Kernel,
};
use std::{collections::HashMap, ffi::c_void, rc::Rc};

#[derive(Debug)]
pub struct RawCL {
    pub ptr: *mut c_void,
    pub host_ptr: *mut u8,
    pub len: usize,
    pub node: Node,
}

impl Drop for RawCL {
    fn drop(&mut self) {
        unsafe { release_mem_object(self.ptr).unwrap() };
    }
}

#[derive(Debug, Default)]
/// This stores the previously compiled OpenCL kernels.
pub struct KernelCacheCL {
    pub kernel_cache: HashMap<String, Rc<Kernel>>,
}

impl KernelCacheCL {
    /// Returns a cached kernel. If the kernel source code does not exist, a new kernel is created and cached.
    ///
    /// # Example
    /// ```
    /// use std::collections::HashMap;
    /// use custos::{OpenCL, opencl::KernelCacheCL};
    ///
    /// fn main() -> custos::Result<()> {
    ///     let device = OpenCL::new(0)?;
    ///     
    ///     let mut kernel_cache = KernelCacheCL {
    ///         kernel_cache: HashMap::new(),
    ///     };
    ///     
    ///     let mut kernel_fn = || kernel_cache.kernel_cache(&device, "
    ///         __kernel void test(__global float* test) {}
    ///     ");
    ///     
    ///     let kernel = kernel_fn()?;
    ///     let same_kernel = kernel_fn()?;
    ///     
    ///     assert_eq!(kernel.0, same_kernel.0);
    ///     Ok(())
    /// }
    /// ```
    pub fn kernel_cache(&mut self, device: &OpenCL, src: &str) -> Result<Rc<Kernel>, Error> {
        if let Some(kernel) = self.kernel_cache.get(src) {
            return Ok(kernel.clone());
        }

        let program = create_program_with_source(&device.ctx(), src)?;
        build_program(&program, &[device.device()], Some("-cl-std=CL1.2"))?; //-cl-single-precision-constant
        let kernel = create_kernels_in_program(&program)?[0].clone();

        self.kernel_cache.insert(src.to_string(), kernel.clone());
        Ok(kernel)
    }
}

#[cfg(test)]
mod tests {
    use super::KernelCacheCL;
    use crate::OpenCL;
    use std::collections::HashMap;

    #[test]
    fn test_kernel_cache() -> crate::Result<()> {
        let device = OpenCL::new(0)?;

        let mut kernel_cache = KernelCacheCL {
            kernel_cache: HashMap::new(),
        };

        let mut kernel_fn = || {
            kernel_cache.kernel_cache(
                &device,
                "
            __kernel void foo(__global float* test) {}
        ",
            )
        };

        let kernel = kernel_fn()?;
        let same_kernel = kernel_fn()?;

        assert_eq!(kernel.0, same_kernel.0);

        let kernel = kernel_fn()?;
        let another_kernel = kernel_cache.kernel_cache(
            &device,
            "
            __kernel void bar(__global float* test, __global float* out) {}
        ",
        )?;

        assert_ne!(kernel.0, another_kernel.0);

        Ok(())
    }
}