unmtx_gpu/
opencl.rs

1//
2// Copyright (c) 2025-2026 Łukasz Szpakowski
3//
4// This Source Code Form is subject to the terms of the Mozilla Public
5// License, v. 2.0. If a copy of the MPL was not distributed with this
6// file, You can obtain one at https://mozilla.org/MPL/2.0/.
7//
8//! A module that contains an OpenCL backend.
9use std::mem::size_of;
10use std::ptr::null_mut;
11use std::sync::Arc;
12use std::sync::Mutex;
13use crate::Backend;
14use crate::BackendArray;
15use crate::Error;
16use crate::Result;
17use crate::mutex_lock;
18
19pub use opencl3::context::Context;
20pub use opencl3::device::Device;
21pub use opencl3::device::CL_DEVICE_TYPE_ACCELERATOR;
22pub use opencl3::device::CL_DEVICE_TYPE_ALL;
23pub use opencl3::device::CL_DEVICE_TYPE_CPU;
24pub use opencl3::device::CL_DEVICE_TYPE_CUSTOM;
25pub use opencl3::device::CL_DEVICE_TYPE_DEFAULT;
26pub use opencl3::device::CL_DEVICE_TYPE_GPU;
27pub use opencl3::device::cl_device_id;
28pub use opencl3::error_codes::ClError;
29pub use opencl3::platform::Platform;
30pub use opencl3::platform::get_platforms;
31
32use cl3::info_type::InfoType;
33use opencl3::command_queue::CommandQueue;
34use opencl3::device::CL_DEVICE_MAX_WORK_GROUP_SIZE;
35use opencl3::device::get_device_info;
36use opencl3::event::Event;
37use opencl3::kernel::ExecuteKernel;
38use opencl3::kernel::Kernel;
39use opencl3::memory::Buffer;
40use opencl3::memory::ClMem;
41use opencl3::memory::cl_mem;
42use opencl3::memory::CL_MEM_READ_WRITE;
43use opencl3::program::Program;
44use opencl3::types::CL_TRUE;
45
46const SOURCE: &'static str = include_str!("opencl.cl");
47
48/// A structure of OpenCL backend array.
49///
50/// This structure contains the reference to an OpenCL buffer.
51#[derive(Debug)]
52pub struct ClBackendArray
53{
54    buffer: Arc<Mutex<Buffer<f32>>>,
55    len: usize,
56}
57
58struct ClInnerBackend
59{
60    context: Context,
61    command_queue: CommandQueue,
62    program: Program,
63    group_size_for_1d: usize,
64    group_size_for_2d: usize,
65}
66
67/// A structure of OpenCL backend.
68pub struct ClBackend
69{
70    inner: Mutex<ClInnerBackend>,
71}
72
73fn preferred_work_sizes(n: usize, m: usize, group_size_for_1d: usize, group_size_for_2d: usize, item_row_count: usize, item_col_count: usize, are_swapped_dims: bool, is_mul: bool) -> (usize, usize, usize, usize)
74{
75    if m <= item_col_count && !is_mul {
76        let n2 = (((n + item_row_count - 1) / item_row_count + group_size_for_1d - 1) / group_size_for_1d) * group_size_for_1d;
77        if !are_swapped_dims {
78            (group_size_for_1d, 1, n2, 1)
79        } else {
80            (1, group_size_for_1d, 1, n2)
81        }
82    } else if n <= item_row_count && !is_mul {
83        let m2 = (((m + item_col_count - 1) / item_col_count + group_size_for_1d - 1) / group_size_for_1d) * group_size_for_1d;
84        if !are_swapped_dims {
85            (1, group_size_for_1d, 1, m2)
86        } else {
87            (group_size_for_1d, 1, m2, 1)
88        }
89    } else if is_mul {
90        let n2 = (((n + 7) / 8 + ((group_size_for_2d + 1) / 2) - 1) / ((group_size_for_2d + 1) / 2)) * ((group_size_for_2d + 1) / 2);
91        let m2 = (((m + 3) / 4 + ((group_size_for_2d + 1) / 2) - 1) / ((group_size_for_2d + 1) / 2)) * ((group_size_for_2d + 1) / 2);
92        if !are_swapped_dims {
93            ((group_size_for_2d + 1) / 2, (group_size_for_2d + 1) / 2, n2, m2)
94        } else {
95            ((group_size_for_2d + 1) / 2, (group_size_for_2d + 1) / 2, m2, n2)
96        }
97    } else {
98        let n2 = (((n + item_row_count - 1) / item_row_count + group_size_for_2d - 1) / group_size_for_2d) * group_size_for_2d;
99        let m2 = (((m + item_col_count - 1) / item_col_count + group_size_for_2d - 1) / group_size_for_2d) * group_size_for_2d;
100        if !are_swapped_dims {
101            (group_size_for_2d, group_size_for_2d, n2, m2)
102        } else {
103            (group_size_for_2d, group_size_for_2d, m2, n2)
104        }
105    }
106}
107
108impl ClBackend
109{
110    /// Creates an OpenCL backend for a first OpenCL platform and a first OpenCL device.
111    pub fn new() -> Result<ClBackend>
112    {
113        let platforms = match get_platforms() {
114            Ok(tmp_platforms) => tmp_platforms,
115            Err(err) => return Err(Error::OpenCl(err)),
116        };
117        if platforms.is_empty() {
118            return Err(Error::NoPlatform);
119        }
120        let device_ids = match platforms[0].get_devices(CL_DEVICE_TYPE_DEFAULT) {
121            Ok(tmp_device_ids) => tmp_device_ids,
122            Err(err) => return Err(Error::OpenCl(err)),
123        };
124        if device_ids.is_empty() {
125            return Err(Error::NoDevice);
126        }
127        let device = Device::new(device_ids[0]);
128        let context = match Context::from_device(&device) {
129            Ok(tmp_context) => tmp_context,
130            Err(err) => return Err(Error::OpenCl(err)),
131        };
132        Self::new_with_context(context)
133    }
134    
135    /// Creates an OpenCL backend with the context.
136    pub fn new_with_context(context: Context) -> Result<ClBackend>
137    {
138        let command_queue = match CommandQueue::create_default_with_properties(&context, 0, 0) {
139            Ok(tmp_command_queue) => tmp_command_queue,
140            Err(err) => return Err(Error::OpenCl(err)),
141        };
142        let program = match Program::create_and_build_from_source(&context, SOURCE, "") {
143            Ok(tmp_program) => tmp_program,
144            Err(msg) => return Err(Error::Compilation(msg)),
145        };
146        let group_size_for_1d = match get_device_info(context.default_device(), CL_DEVICE_MAX_WORK_GROUP_SIZE) {
147            Ok(InfoType::Size(tmp_group_size_for_1d)) => tmp_group_size_for_1d,
148            _ => return Err(Error::InvalidDeviceInfoType),
149        };
150        let group_size_for_2d = (group_size_for_1d as f64).sqrt().floor() as usize;
151        let inner = ClInnerBackend {
152            context,
153            command_queue,
154            program,
155            group_size_for_1d,
156            group_size_for_2d,
157        };
158        Ok(ClBackend { inner: Mutex::new(inner), })
159    }
160    
161    fn check_and_enqueue_nd_range2<F, G>(&self, kernel_name: &str, a: &BackendArray, b: &BackendArray, f: F, g: G) -> Result<()>
162        where F: FnOnce(&ClBackendArray, &ClBackendArray) -> Result<()>,
163            G: FnOnce(&ClInnerBackend, &Kernel, cl_mem, cl_mem) -> Result<Event>
164    {
165        #[allow(unreachable_patterns)]
166        match (a, b) {
167            (BackendArray::OpenCl(a2), BackendArray::OpenCl(b2)) => {
168                f(a2, b2)?;
169                let inner_g = mutex_lock(&self.inner)?;
170                let kernel = match Kernel::create(&inner_g.program, kernel_name) {
171                    Ok(tmp_kernel) => tmp_kernel,
172                    Err(err) => return Err(Error::OpenCl(err)),
173                };
174                let event = if !Arc::ptr_eq(&a2.buffer, &b2.buffer) {
175                    let a_buffer_g = mutex_lock(&a2.buffer)?;
176                    let mut b_buffer_g = mutex_lock(&b2.buffer)?;
177                    g(&*inner_g, &kernel, a_buffer_g.get(), b_buffer_g.get_mut())?
178                } else {
179                    let mut a_buffer_g = mutex_lock(&a2.buffer)?;
180                    g(&*inner_g, &kernel, a_buffer_g.get(), a_buffer_g.get_mut())?
181                };
182                match event.wait() {
183                    Ok(()) => (),
184                    Err(err) => return Err(Error::OpenCl(err)),
185                }
186                Ok(())
187            },
188            _ => Err(Error::InvalidBackendArray),
189        }
190    }
191
192    fn check_and_enqueue_nd_range3<F, G>(&self, kernel_name: &str, a: &BackendArray, b: &BackendArray, c: &BackendArray, f: F, g: G) -> Result<()>
193        where F: FnOnce(&ClBackendArray, &ClBackendArray, &ClBackendArray) -> Result<()>,
194            G: FnOnce(&ClInnerBackend, &Kernel, cl_mem, cl_mem, cl_mem) -> Result<Event>
195    {
196        #[allow(unreachable_patterns)]
197        match (a, b, c) {
198            (BackendArray::OpenCl(a2), BackendArray::OpenCl(b2), BackendArray::OpenCl(c2)) => {
199                f(a2, b2, c2)?;
200                let inner_g = mutex_lock(&self.inner)?;
201                let kernel = match Kernel::create(&inner_g.program, kernel_name) {
202                    Ok(tmp_kernel) => tmp_kernel,
203                    Err(err) => return Err(Error::OpenCl(err)),
204                };
205                let event = match (Arc::ptr_eq(&a2.buffer, &b2.buffer), Arc::ptr_eq(&a2.buffer, &c2.buffer), Arc::ptr_eq(&b2.buffer, &c2.buffer)) {
206                    (false, false, false) => {
207                        let a_buffer_g = mutex_lock(&a2.buffer)?;
208                        let b_buffer_g = mutex_lock(&b2.buffer)?;
209                        let mut c_buffer_g = mutex_lock(&c2.buffer)?;
210                        g(&*inner_g, &kernel, a_buffer_g.get(), b_buffer_g.get(), c_buffer_g.get_mut())?
211                    },
212                    (true, false, false) => {
213                        let a_buffer_g = mutex_lock(&a2.buffer)?;
214                        let mut c_buffer_g = mutex_lock(&c2.buffer)?;
215                        g(&*inner_g, &kernel, a_buffer_g.get(), a_buffer_g.get(), c_buffer_g.get_mut())?
216                    },
217                    (false, true, false) => {
218                        let mut a_buffer_g = mutex_lock(&a2.buffer)?;
219                        let b_buffer_g = mutex_lock(&b2.buffer)?;
220                        g(&*inner_g, &kernel, a_buffer_g.get(), b_buffer_g.get(), a_buffer_g.get_mut())?
221                    },
222                    (false, false, true) => {
223                        let a_buffer_g = mutex_lock(&a2.buffer)?;
224                        let mut b_buffer_g = mutex_lock(&b2.buffer)?;
225                        g(&*inner_g, &kernel, a_buffer_g.get(), b_buffer_g.get(), b_buffer_g.get_mut())?
226                    },
227                    _ => {
228                        let mut a_buffer_g = mutex_lock(&a2.buffer)?;
229                        g(&*inner_g, &kernel, a_buffer_g.get(), a_buffer_g.get(), a_buffer_g.get_mut())?
230                    },
231                };
232                match event.wait() {
233                    Ok(()) => (),
234                    Err(err) => return Err(Error::OpenCl(err)),
235                }
236                Ok(())
237            },
238            _ => Err(Error::InvalidBackendArray),
239        }
240    }
241    
242    fn check_and_enqueue_nd_range_for_fun(&self, kernel_name: &str, a: &BackendArray, b: &BackendArray, n: usize, m: usize, item_row_count: usize, item_col_count: usize, are_swapped_dims: bool) -> Result<()>
243    {
244        self.check_and_enqueue_nd_range2(kernel_name, a, b, |a2, b2| {
245                if a2.len != n * m {
246                    return Err(Error::BackendArrayElemCount(a2.len, n * m));
247                }
248                if b2.len != n * m {
249                    return Err(Error::BackendArrayElemCount(b2.len, n * m));
250                }
251                Ok(())
252        }, |inner, kernel, a_mem, b_mem| {
253                let n2 = n as u64;
254                let m2 = m as u64;
255                let (n3, m3, n4, m4) = preferred_work_sizes(n, m, inner.group_size_for_1d, inner.group_size_for_2d, item_row_count, item_col_count, are_swapped_dims, false);
256                unsafe {
257                    let res = ExecuteKernel::new(kernel)
258                    .set_arg(&a_mem)
259                    .set_arg(&b_mem)
260                    .set_arg(&n2)
261                    .set_arg(&m2)
262                    .set_local_work_sizes(&[n3, m3])
263                    .set_global_work_sizes(&[n4, m4])
264                    .enqueue_nd_range(&inner.command_queue);
265                    match res {
266                        Ok(event) => Ok(event),
267                        Err(err) => Err(Error::OpenCl(err)),
268                    }
269                }
270        })
271    }
272
273    fn check_and_enqueue_nd_range_for_op(&self, kernel_name: &str, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize, item_row_count: usize, item_col_count: usize, are_swapped_dims: bool) -> Result<()>
274    {
275        self.check_and_enqueue_nd_range3(kernel_name, a, b, c, |a2, b2, c2| {
276                if a2.len != n * m {
277                    return Err(Error::BackendArrayElemCount(a2.len, n * m));
278                }
279                if b2.len != n * m {
280                    return Err(Error::BackendArrayElemCount(b2.len, n * m));
281                }
282                if c2.len != n * m {
283                    return Err(Error::BackendArrayElemCount(c2.len, n * m));
284                }
285                Ok(())
286        }, |inner, kernel, a_mem, b_mem, c_mem| {
287                let n2 = n as u64;
288                let m2 = m as u64;
289                let (n3, m3, n4, m4) = preferred_work_sizes(n, m, inner.group_size_for_1d, inner.group_size_for_2d, item_row_count, item_col_count, are_swapped_dims, false);
290                unsafe {
291                    let res = ExecuteKernel::new(kernel)
292                    .set_arg(&a_mem)
293                    .set_arg(&b_mem)
294                    .set_arg(&c_mem)
295                    .set_arg(&n2)
296                    .set_arg(&m2)
297                    .set_local_work_sizes(&[n3, m3])
298                    .set_global_work_sizes(&[n4, m4])
299                    .enqueue_nd_range(&inner.command_queue);
300                    match res {
301                        Ok(event) => Ok(event),
302                        Err(err) => Err(Error::OpenCl(err)),
303                    }
304                }
305        })
306    }
307
308    fn check_and_enqueue_nd_range_for_mul(&self, kernel_name: &str, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize, l: usize, item_row_count: usize, item_col_count: usize, are_swapped_dims: bool) -> Result<()>
309    {
310        self.check_and_enqueue_nd_range3(kernel_name, a, b, c, |a2, b2, c2| {
311                if a2.len != n * l {
312                    return Err(Error::BackendArrayElemCount(a2.len, n * l));
313                }
314                if b2.len != l * m {
315                    return Err(Error::BackendArrayElemCount(b2.len, l * m));
316                }
317                if c2.len != n * m {
318                    return Err(Error::BackendArrayElemCount(c2.len, n * m));
319                }
320                Ok(())
321        }, |inner, kernel, a_mem, b_mem, c_mem| {
322                let n2 = n as u64;
323                let m2 = m as u64;
324                let l2 = l as u64;
325                let (n3, m3, n4, m4) = preferred_work_sizes(n, m, inner.group_size_for_1d, inner.group_size_for_2d, item_row_count, item_col_count, are_swapped_dims, true);
326                unsafe {
327                    let res = ExecuteKernel::new(kernel)
328                    .set_arg(&a_mem)
329                    .set_arg(&b_mem)
330                    .set_arg(&c_mem)
331                    .set_arg_local_buffer(n3 * m3 * 8 * size_of::<f32>())
332                    .set_arg_local_buffer(n3 * m3 * 4 * size_of::<f32>())
333                    .set_arg(&n2)
334                    .set_arg(&m2)
335                    .set_arg(&l2)
336                    .set_local_work_sizes(&[n3, m3])
337                    .set_global_work_sizes(&[n4, m4])
338                    .enqueue_nd_range(&inner.command_queue);
339                    match res {
340                        Ok(event) => Ok(event),
341                        Err(err) => Err(Error::OpenCl(err)),
342                    }
343                }
344        })
345    }
346
347    fn check_and_enqueue_nd_range_for_scalar(&self, kernel_name: &str, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize, item_row_count: usize, item_col_count: usize, are_swapped_dims: bool) -> Result<()>
348    {
349        self.check_and_enqueue_nd_range2(kernel_name, a, c, |a2, c2| {
350                if a2.len != n * m  {
351                    return Err(Error::BackendArrayElemCount(a2.len, n * m));
352                }
353                if c2.len != n * m {
354                    return Err(Error::BackendArrayElemCount(c2.len, n * m));
355                }
356                Ok(())
357        }, |inner, kernel, a_mem, c_mem| {
358                let n2 = n as u64;
359                let m2 = m as u64;
360                let (n3, m3, n4, m4) = preferred_work_sizes(n, m, inner.group_size_for_1d, inner.group_size_for_2d, item_row_count, item_col_count, are_swapped_dims, false);
361                unsafe {
362                    let res = ExecuteKernel::new(kernel)
363                    .set_arg(&a_mem)
364                    .set_arg(&b)
365                    .set_arg(&c_mem)
366                    .set_arg(&n2)
367                    .set_arg(&m2)
368                    .set_local_work_sizes(&[n3, m3])
369                    .set_global_work_sizes(&[n4, m4])
370                    .enqueue_nd_range(&inner.command_queue);
371                    match res {
372                        Ok(event) => Ok(event),
373                        Err(err) => Err(Error::OpenCl(err)),
374                    }
375                }
376        })
377    }
378
379    fn check_and_enqueue_nd_range_for_fun_and_tiles(&self, kernel_name: &str, a: &BackendArray, b: &BackendArray, n: usize, m: usize, item_row_count: usize, item_col_count: usize, are_swapped_dims: bool) -> Result<()>
380    {
381        self.check_and_enqueue_nd_range2(kernel_name, a, b, |a2, b2| {
382                if a2.len != n * m {
383                    return Err(Error::BackendArrayElemCount(a2.len, n * m));
384                }
385                if b2.len != n * m {
386                    return Err(Error::BackendArrayElemCount(b2.len, n * m));
387                }
388                Ok(())
389        }, |inner, kernel, a_mem, b_mem| {
390                let n2 = n as u64;
391                let m2 = m as u64;
392                let (n3, m3, n4, m4) = preferred_work_sizes(n, m, inner.group_size_for_1d, inner.group_size_for_2d, item_row_count, item_col_count, are_swapped_dims, false);
393                unsafe {
394                    let res = ExecuteKernel::new(kernel)
395                    .set_arg(&a_mem)
396                    .set_arg(&b_mem)
397                    .set_arg_local_buffer(n3 * item_row_count * m3 * item_col_count * size_of::<f32>())
398                    .set_arg(&n2)
399                    .set_arg(&m2)
400                    .set_local_work_sizes(&[n3, m3])
401                    .set_global_work_sizes(&[n4, m4])
402                    .enqueue_nd_range(&inner.command_queue);
403                    match res {
404                        Ok(event) => Ok(event),
405                        Err(err) => Err(Error::OpenCl(err)),
406                    }
407                }
408        })
409    }
410
411    fn check_and_enqueue_nd_range_for_repeat_col(&self, kernel_name: &str, a: &BackendArray, b: &BackendArray, n: usize, m: usize, item_row_count: usize, item_col_count: usize, are_swapped_dims: bool) -> Result<()>
412    {
413        self.check_and_enqueue_nd_range2(kernel_name, a, b, |a2, b2| {
414                if a2.len != n {
415                    return Err(Error::BackendArrayElemCount(a2.len, n));
416                }
417                if b2.len != n * m {
418                    return Err(Error::BackendArrayElemCount(b2.len, n * m));
419                }
420                Ok(())
421        }, |inner, kernel, a_mem, b_mem| {
422                let n2 = n as u64;
423                let m2 = m as u64;
424                let (n3, m3, n4, m4) = preferred_work_sizes(n, m, inner.group_size_for_1d, inner.group_size_for_2d, item_row_count, item_col_count, are_swapped_dims, false);
425                unsafe {
426                    let res = ExecuteKernel::new(kernel)
427                    .set_arg(&a_mem)
428                    .set_arg(&b_mem)
429                    .set_arg(&n2)
430                    .set_arg(&m2)
431                    .set_local_work_sizes(&[n3, m3])
432                    .set_global_work_sizes(&[n4, m4])
433                    .enqueue_nd_range(&inner.command_queue);
434                    match res {
435                        Ok(event) => Ok(event),
436                        Err(err) => Err(Error::OpenCl(err)),
437                    }
438                }
439        })
440    }
441
442    fn check_and_enqueue_nd_range_for_repeat_row(&self, kernel_name: &str, a: &BackendArray, b: &BackendArray, n: usize, m: usize, item_row_count: usize, item_col_count: usize, are_swapped_dims: bool) -> Result<()>
443    {
444        self.check_and_enqueue_nd_range2(kernel_name, a, b, |a2, b2| {
445                if a2.len != m {
446                    return Err(Error::BackendArrayElemCount(a2.len, m));
447                }
448                if b2.len != n * m {
449                    return Err(Error::BackendArrayElemCount(b2.len, n * m));
450                }
451                Ok(())
452        }, |inner, kernel, a_mem, b_mem| {
453                let n2 = n as u64;
454                let m2 = m as u64;
455                let (n3, m3, n4, m4) = preferred_work_sizes(n, m, inner.group_size_for_1d, inner.group_size_for_2d, item_row_count, item_col_count, are_swapped_dims, false);
456                unsafe {
457                    let res = ExecuteKernel::new(kernel)
458                    .set_arg(&a_mem)
459                    .set_arg(&b_mem)
460                    .set_arg(&n2)
461                    .set_arg(&m2)
462                    .set_local_work_sizes(&[n3, m3])
463                    .set_global_work_sizes(&[n4, m4])
464                    .enqueue_nd_range(&inner.command_queue);
465                    match res {
466                        Ok(event) => Ok(event),
467                        Err(err) => Err(Error::OpenCl(err)),
468                    }
469                }
470        })
471    }
472}
473
474impl Backend for ClBackend
475{
476    fn name(&self) -> &'static str
477    { "OpenCL" }
478    
479    fn has_cublas(&self) -> bool
480    { false }
481    
482    unsafe fn alloc(&self, n: usize) -> Result<BackendArray>
483    {
484        let inner_g = mutex_lock(&self.inner)?;
485        let buffer: Buffer<f32> = match Buffer::create(&inner_g.context, CL_MEM_READ_WRITE, n, null_mut()) {
486            Ok(tmp_buffer) => tmp_buffer,
487            Err(err) => return Err(Error::OpenCl(err)),
488        };
489        let cl_array = ClBackendArray { buffer: Arc::new(Mutex::new(buffer)), len: n, };
490        Ok(BackendArray::OpenCl(cl_array))
491    }
492
493    fn alloc_and_store_zeros(&self, n: usize) -> Result<BackendArray>
494    {
495        let inner_g = mutex_lock(&self.inner)?;
496        let mut buffer: Buffer<f32> = match unsafe { Buffer::create(&inner_g.context, CL_MEM_READ_WRITE, n, null_mut()) } {
497            Ok(tmp_buffer) => tmp_buffer,
498            Err(err) => return Err(Error::OpenCl(err)),
499        };
500        let event = match unsafe { inner_g.command_queue.enqueue_fill_buffer(&mut buffer, &[0.0f32], 0, n * size_of::<f32>(), &[]) } {
501            Ok(tmp_event) => tmp_event,
502            Err(err) => return Err(Error::OpenCl(err)),
503        };
504        match event.wait() {
505            Ok(()) => (),
506            Err(err) => return Err(Error::OpenCl(err)),
507        }
508        let cl_array = ClBackendArray { buffer: Arc::new(Mutex::new(buffer)), len: n, };
509        Ok(BackendArray::OpenCl(cl_array))
510    }
511    
512    fn alloc_and_store(&self, elems: &[f32]) -> Result<BackendArray>
513    {
514        let inner_g = mutex_lock(&self.inner)?;
515        let mut buffer: Buffer<f32> = match unsafe { Buffer::create(&inner_g.context, CL_MEM_READ_WRITE, elems.len(), null_mut()) } {
516            Ok(tmp_buffer) => tmp_buffer,
517            Err(err) => return Err(Error::OpenCl(err)),
518        };
519        let event = match unsafe { inner_g.command_queue.enqueue_write_buffer(&mut buffer, CL_TRUE, 0, elems, &[]) } {
520            Ok(tmp_event) => tmp_event,
521            Err(err) => return Err(Error::OpenCl(err)),
522        };
523        match event.wait() {
524            Ok(()) => (),
525            Err(err) => return Err(Error::OpenCl(err)),
526        }
527        let cl_array = ClBackendArray { buffer: Arc::new(Mutex::new(buffer)), len: elems.len(), };
528        Ok(BackendArray::OpenCl(cl_array))
529    }
530    
531    fn load(&self, a: &BackendArray, elems: &mut [f32]) -> Result<()>
532    {
533        #[allow(unreachable_patterns)]
534        match a {
535            BackendArray::OpenCl(a2) => {
536                if a2.len != elems.len() {
537                    return Err(Error::BackendArrayElemCount(a2.len, elems.len()));
538                }
539                let inner_g = mutex_lock(&self.inner)?;
540                let a_buffer_g = mutex_lock(&a2.buffer)?;
541                let event = match unsafe { inner_g.command_queue.enqueue_read_buffer(&*a_buffer_g, CL_TRUE, 0, elems, &[]) } {
542                    Ok(tmp_event) => tmp_event,
543                    Err(err) => return Err(Error::OpenCl(err)),
544                };
545                match event.wait() {
546                    Ok(()) => (),
547                    Err(err) => return Err(Error::OpenCl(err)),
548                }
549            },
550            _ => return Err(Error::InvalidBackendArray),
551        }
552        Ok(())
553    }
554
555    fn store(&self, a: &BackendArray, elems: &[f32]) -> Result<()>
556    {
557        #[allow(unreachable_patterns)]
558        match a {
559            BackendArray::OpenCl(a2) => {
560                if a2.len != elems.len() {
561                    return Err(Error::BackendArrayElemCount(a2.len, elems.len()));
562                }
563                let inner_g = mutex_lock(&self.inner)?;
564                let mut a_buffer_g = mutex_lock(&a2.buffer)?;
565                let event = match unsafe { inner_g.command_queue.enqueue_write_buffer(&mut *a_buffer_g, CL_TRUE, 0, elems, &[]) } {
566                    Ok(tmp_event) => tmp_event,
567                    Err(err) => return Err(Error::OpenCl(err)),
568                };
569                match event.wait() {
570                    Ok(()) => (),
571                    Err(err) => return Err(Error::OpenCl(err)),
572                }
573            },
574            _ => return Err(Error::InvalidBackendArray),
575        }
576        Ok(())
577    }
578    
579    fn copy(&self, a: &BackendArray, b: &BackendArray) -> Result<()>
580    {
581        #[allow(unreachable_patterns)]
582        match (a, b) {
583            (BackendArray::OpenCl(a2), BackendArray::OpenCl(b2)) => {
584                if Arc::ptr_eq(&a2.buffer, &b2.buffer) {
585                    return Ok(());
586                }
587                if a2.len != b2.len {
588                    return Err(Error::TwoBackendArrayElemCounts(a2.len, b2.len));
589                }
590                let inner_g = mutex_lock(&self.inner)?;
591                let a_buffer_g = mutex_lock(&a2.buffer)?;
592                let mut b_buffer_g = mutex_lock(&b2.buffer)?;
593                let event = match unsafe { inner_g.command_queue.enqueue_copy_buffer(&*a_buffer_g, &mut *b_buffer_g, 0, 0, a2.len * size_of::<f32>(), &[]) } {
594                    Ok(tmp_event) => tmp_event,
595                    Err(err) => return Err(Error::OpenCl(err)),
596                };
597                match event.wait() {
598                    Ok(()) => (),
599                    Err(err) => return Err(Error::OpenCl(err)),
600                }
601            },
602            _ => return Err(Error::InvalidBackendArray),
603        }
604        Ok(())
605    }
606
607    fn transpose_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
608    { self.check_and_enqueue_nd_range_for_fun("transpose_a", a, b, n, m, 2, 2, true) }
609
610    fn add_a_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
611    { self.check_and_enqueue_nd_range_for_op("add_a_b", a, b, c, n, m, 2, 2, true) }
612
613    fn add_at_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
614    { self.check_and_enqueue_nd_range_for_op("add_at_b", a, b, c, n, m, 2, 2, true) }
615    
616    fn add_a_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
617    { self.check_and_enqueue_nd_range_for_op("add_a_bt", a, b, c, n, m, 2, 2, true) }
618
619    fn add_at_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
620    { self.check_and_enqueue_nd_range_for_op("add_at_bt", a, b, c, n, m, 2, 2, true) }
621
622    fn sub_a_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
623    { self.check_and_enqueue_nd_range_for_op("sub_a_b", a, b, c, n, m, 2, 2, true) }
624
625    fn sub_at_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
626    { self.check_and_enqueue_nd_range_for_op("sub_at_b", a, b, c, n, m, 2, 2, true) }
627    
628    fn sub_a_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
629    { self.check_and_enqueue_nd_range_for_op("sub_a_bt", a, b, c, n, m, 2, 2, true) }
630
631    fn sub_at_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>    
632    { self.check_and_enqueue_nd_range_for_op("sub_at_bt", a, b, c, n, m, 2, 2, true) }
633    
634    fn mul_a_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize, l: usize) -> Result<()>
635    { self.check_and_enqueue_nd_range_for_mul("mul_a_b", a, b, c, n, m, l, 8, 4, true) }
636
637    fn mul_at_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize, l: usize) -> Result<()>
638    { self.check_and_enqueue_nd_range_for_mul("mul_at_b", a, b, c, n, m, l, 8, 4, false) }
639
640    fn mul_a_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize, l: usize) -> Result<()>
641    { self.check_and_enqueue_nd_range_for_mul("mul_a_bt", a, b, c, n, m, l, 8, 4, true) }
642
643    fn mul_at_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize, l: usize) -> Result<()>
644    { self.check_and_enqueue_nd_range_for_mul("mul_at_bt", a, b, c, n, m, l, 8, 4, false) }
645
646    fn mul_a_b_for_elems(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
647    { self.check_and_enqueue_nd_range_for_op("mul_a_b_for_elems", a, b, c, n, m, 2, 2, true) }
648
649    fn mul_at_b_for_elems(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
650    { self.check_and_enqueue_nd_range_for_op("mul_at_b_for_elems", a, b, c, n, m, 2, 2, true) }
651    
652    fn mul_a_bt_for_elems(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
653    { self.check_and_enqueue_nd_range_for_op("mul_a_bt_for_elems", a, b, c, n, m, 2, 2, true) }
654    
655    fn mul_at_bt_for_elems(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
656    { self.check_and_enqueue_nd_range_for_op("mul_at_bt_for_elems", a, b, c, n, m, 2, 2, true) }
657
658    fn div_a_b_for_elems(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
659    { self.check_and_enqueue_nd_range_for_op("div_a_b_for_elems", a, b, c, n, m, 2, 2, true) }
660
661    fn div_at_b_for_elems(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
662    { self.check_and_enqueue_nd_range_for_op("div_at_b_for_elems", a, b, c, n, m, 2, 2, true) }
663    
664    fn div_a_bt_for_elems(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
665    { self.check_and_enqueue_nd_range_for_op("div_a_bt_for_elems", a, b, c, n, m, 2, 2, true) }
666    
667    fn div_at_bt_for_elems(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
668    { self.check_and_enqueue_nd_range_for_op("div_at_bt_for_elems", a, b, c, n, m, 2, 2, true) }
669
670    fn add_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
671    { self.check_and_enqueue_nd_range_for_scalar("add_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
672
673    fn add_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
674    { self.check_and_enqueue_nd_range_for_scalar("add_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
675
676    fn sub_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
677    { self.check_and_enqueue_nd_range_for_scalar("sub_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
678
679    fn sub_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
680    { self.check_and_enqueue_nd_range_for_scalar("sub_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
681
682    fn rsub_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
683    { self.check_and_enqueue_nd_range_for_scalar("rsub_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
684
685    fn rsub_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
686    { self.check_and_enqueue_nd_range_for_scalar("rsub_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
687    
688    fn mul_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
689    { self.check_and_enqueue_nd_range_for_scalar("mul_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
690
691    fn mul_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
692    { self.check_and_enqueue_nd_range_for_scalar("mul_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
693
694    fn div_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
695    { self.check_and_enqueue_nd_range_for_scalar("div_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
696
697    fn div_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
698    { self.check_and_enqueue_nd_range_for_scalar("div_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
699
700    fn rdiv_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
701    { self.check_and_enqueue_nd_range_for_scalar("rdiv_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
702
703    fn rdiv_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
704    { self.check_and_enqueue_nd_range_for_scalar("rdiv_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
705
706    fn sigmoid_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
707    { self.check_and_enqueue_nd_range_for_fun("sigmoid_a", a, b, n, m, 2, 2, true) }
708
709    fn sigmoid_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
710    { self.check_and_enqueue_nd_range_for_fun("sigmoid_at", a, b, n, m, 2, 2, true) }
711
712    fn tanh_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
713    { self.check_and_enqueue_nd_range_for_fun("tanh_a", a, b, n, m, 2, 2, true) }
714
715    fn tanh_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
716    { self.check_and_enqueue_nd_range_for_fun("tanh_at", a, b, n, m, 2, 2, true) }
717
718    fn swish_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
719    { self.check_and_enqueue_nd_range_for_fun("swish_a", a, b, n, m, 2, 2, true) }
720
721    fn swish_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
722    { self.check_and_enqueue_nd_range_for_fun("swish_at", a, b, n, m, 2, 2, true) }
723
724    fn softmax_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
725    { self.check_and_enqueue_nd_range_for_fun_and_tiles("softmax_a", a, b, n, m, 2, 2, true) }
726
727    fn softmax_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
728    { self.check_and_enqueue_nd_range_for_fun_and_tiles("softmax_at", a, b, n, m, 2, 2, false) }
729
730    fn sqrt_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
731    { self.check_and_enqueue_nd_range_for_fun("sqrt_a", a, b, n, m, 2, 2, true) }
732
733    fn sqrt_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
734    { self.check_and_enqueue_nd_range_for_fun("sqrt_at", a, b, n, m, 2, 2, true) }
735    
736    fn repeat_col_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
737    { self.check_and_enqueue_nd_range_for_repeat_col("repeat_col_a", a, b, n, m, 2, 2, true) }
738
739    fn repeat_row_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
740    { self.check_and_enqueue_nd_range_for_repeat_row("repeat_row_a", a, b, n, m, 2, 2, true) }
741
742    fn abs_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
743    { self.check_and_enqueue_nd_range_for_fun("abs_a", a, b, n, m, 2, 2, true) }
744
745    fn abs_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
746    { self.check_and_enqueue_nd_range_for_fun("abs_at", a, b, n, m, 2, 2, true) }
747
748    fn pow_a_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
749    { self.check_and_enqueue_nd_range_for_op("pow_a_b", a, b, c, n, m, 2, 2, true) }
750
751    fn pow_at_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
752    { self.check_and_enqueue_nd_range_for_op("pow_at_b", a, b, c, n, m, 2, 2, true) }
753    
754    fn pow_a_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
755    { self.check_and_enqueue_nd_range_for_op("pow_a_bt", a, b, c, n, m, 2, 2, true) }
756    
757    fn pow_at_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
758    { self.check_and_enqueue_nd_range_for_op("pow_at_bt", a, b, c, n, m, 2, 2, true) }
759
760    fn pow_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
761    { self.check_and_enqueue_nd_range_for_scalar("pow_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
762
763    fn pow_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
764    { self.check_and_enqueue_nd_range_for_scalar("pow_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
765
766    fn rpow_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
767    { self.check_and_enqueue_nd_range_for_scalar("rpow_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
768
769    fn rpow_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
770    { self.check_and_enqueue_nd_range_for_scalar("rpow_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
771
772    fn exp_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
773    { self.check_and_enqueue_nd_range_for_fun("exp_a", a, b, n, m, 2, 2, true) }
774
775    fn exp_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
776    { self.check_and_enqueue_nd_range_for_fun("exp_at", a, b, n, m, 2, 2, true) }
777
778    fn ln_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
779    { self.check_and_enqueue_nd_range_for_fun("ln_a", a, b, n, m, 2, 2, true) }
780
781    fn ln_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
782    { self.check_and_enqueue_nd_range_for_fun("ln_at", a, b, n, m, 2, 2, true) }
783
784    fn log2_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
785    { self.check_and_enqueue_nd_range_for_fun("log2_a", a, b, n, m, 2, 2, true) }
786
787    fn log2_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
788    { self.check_and_enqueue_nd_range_for_fun("log2_at", a, b, n, m, 2, 2, true) }
789
790    fn log10_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
791    { self.check_and_enqueue_nd_range_for_fun("log10_a", a, b, n, m, 2, 2, true) }
792
793    fn log10_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
794    { self.check_and_enqueue_nd_range_for_fun("log10_at", a, b, n, m, 2, 2, true) }
795
796    fn sin_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
797    { self.check_and_enqueue_nd_range_for_fun("sin_a", a, b, n, m, 2, 2, true) }
798
799    fn sin_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
800    { self.check_and_enqueue_nd_range_for_fun("sin_at", a, b, n, m, 2, 2, true) }
801
802    fn cos_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
803    { self.check_and_enqueue_nd_range_for_fun("cos_a", a, b, n, m, 2, 2, true) }
804
805    fn cos_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
806    { self.check_and_enqueue_nd_range_for_fun("cos_at", a, b, n, m, 2, 2, true) }
807
808    fn tan_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
809    { self.check_and_enqueue_nd_range_for_fun("tan_a", a, b, n, m, 2, 2, true) }
810
811    fn tan_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
812    { self.check_and_enqueue_nd_range_for_fun("tan_at", a, b, n, m, 2, 2, true) }
813
814    fn asin_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
815    { self.check_and_enqueue_nd_range_for_fun("asin_a", a, b, n, m, 2, 2, true) }
816
817    fn asin_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
818    { self.check_and_enqueue_nd_range_for_fun("asin_at", a, b, n, m, 2, 2, true) }
819
820    fn acos_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
821    { self.check_and_enqueue_nd_range_for_fun("acos_a", a, b, n, m, 2, 2, true) }
822
823    fn acos_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
824    { self.check_and_enqueue_nd_range_for_fun("acos_at", a, b, n, m, 2, 2, true) }
825
826    fn atan_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
827    { self.check_and_enqueue_nd_range_for_fun("atan_a", a, b, n, m, 2, 2, true) }
828
829    fn atan_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
830    { self.check_and_enqueue_nd_range_for_fun("atan_at", a, b, n, m, 2, 2, true) }
831
832    fn atan2_a_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
833    { self.check_and_enqueue_nd_range_for_op("atan2_a_b", a, b, c, n, m, 2, 2, true) }
834
835    fn atan2_at_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
836    { self.check_and_enqueue_nd_range_for_op("atan2_at_b", a, b, c, n, m, 2, 2, true) }
837    
838    fn atan2_a_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
839    { self.check_and_enqueue_nd_range_for_op("atan2_a_bt", a, b, c, n, m, 2, 2, true) }
840    
841    fn atan2_at_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
842    { self.check_and_enqueue_nd_range_for_op("atan2_at_bt", a, b, c, n, m, 2, 2, true) }
843
844    fn atan2_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
845    { self.check_and_enqueue_nd_range_for_scalar("atan2_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
846
847    fn atan2_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
848    { self.check_and_enqueue_nd_range_for_scalar("atan2_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
849
850    fn ratan2_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
851    { self.check_and_enqueue_nd_range_for_scalar("ratan2_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
852
853    fn ratan2_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
854    { self.check_and_enqueue_nd_range_for_scalar("ratan2_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
855
856    fn sinh_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
857    { self.check_and_enqueue_nd_range_for_fun("sinh_a", a, b, n, m, 2, 2, true) }
858
859    fn sinh_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
860    { self.check_and_enqueue_nd_range_for_fun("sinh_at", a, b, n, m, 2, 2, true) }
861
862    fn cosh_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
863    { self.check_and_enqueue_nd_range_for_fun("cosh_a", a, b, n, m, 2, 2, true) }
864
865    fn cosh_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
866    { self.check_and_enqueue_nd_range_for_fun("cosh_at", a, b, n, m, 2, 2, true) }
867
868    fn asinh_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
869    { self.check_and_enqueue_nd_range_for_fun("asinh_a", a, b, n, m, 2, 2, true) }
870
871    fn asinh_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
872    { self.check_and_enqueue_nd_range_for_fun("asinh_at", a, b, n, m, 2, 2, true) }
873
874    fn acosh_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
875    { self.check_and_enqueue_nd_range_for_fun("acosh_a", a, b, n, m, 2, 2, true) }
876
877    fn acosh_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
878    { self.check_and_enqueue_nd_range_for_fun("acosh_at", a, b, n, m, 2, 2, true) }
879
880    fn atanh_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
881    { self.check_and_enqueue_nd_range_for_fun("atanh_a", a, b, n, m, 2, 2, true) }
882
883    fn atanh_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
884    { self.check_and_enqueue_nd_range_for_fun("atanh_at", a, b, n, m, 2, 2, true) }
885
886    fn signum_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
887    { self.check_and_enqueue_nd_range_for_fun("signum_a", a, b, n, m, 2, 2, true) }
888
889    fn signum_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
890    { self.check_and_enqueue_nd_range_for_fun("signum_at", a, b, n, m, 2, 2, true) }
891
892    fn ceil_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
893    { self.check_and_enqueue_nd_range_for_fun("ceil_a", a, b, n, m, 2, 2, true) }
894
895    fn ceil_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
896    { self.check_and_enqueue_nd_range_for_fun("ceil_at", a, b, n, m, 2, 2, true) }
897
898    fn floor_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
899    { self.check_and_enqueue_nd_range_for_fun("floor_a", a, b, n, m, 2, 2, true) }
900
901    fn floor_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
902    { self.check_and_enqueue_nd_range_for_fun("floor_at", a, b, n, m, 2, 2, true) }
903
904    fn round_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
905    { self.check_and_enqueue_nd_range_for_fun("round_a", a, b, n, m, 2, 2, true) }
906
907    fn round_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
908    { self.check_and_enqueue_nd_range_for_fun("round_at", a, b, n, m, 2, 2, true) }
909
910    fn trunc_a(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
911    { self.check_and_enqueue_nd_range_for_fun("trunc_a", a, b, n, m, 2, 2, true) }
912
913    fn trunc_at(&self, a: &BackendArray, b: &BackendArray, n: usize, m: usize) -> Result<()>
914    { self.check_and_enqueue_nd_range_for_fun("trunc_at", a, b, n, m, 2, 2, true) }
915
916    fn max_a_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
917    { self.check_and_enqueue_nd_range_for_op("max_a_b", a, b, c, n, m, 2, 2, true) }
918
919    fn max_at_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
920    { self.check_and_enqueue_nd_range_for_op("max_at_b", a, b, c, n, m, 2, 2, true) }
921    
922    fn max_a_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
923    { self.check_and_enqueue_nd_range_for_op("max_a_bt", a, b, c, n, m, 2, 2, true) }
924    
925    fn max_at_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
926    { self.check_and_enqueue_nd_range_for_op("max_at_bt", a, b, c, n, m, 2, 2, true) }
927
928    fn max_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
929    { self.check_and_enqueue_nd_range_for_scalar("max_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
930
931    fn max_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
932    { self.check_and_enqueue_nd_range_for_scalar("max_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
933
934    fn min_a_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
935    { self.check_and_enqueue_nd_range_for_op("min_a_b", a, b, c, n, m, 2, 2, true) }
936
937    fn min_at_b(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
938    { self.check_and_enqueue_nd_range_for_op("min_at_b", a, b, c, n, m, 2, 2, true) }
939    
940    fn min_a_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
941    { self.check_and_enqueue_nd_range_for_op("min_a_bt", a, b, c, n, m, 2, 2, true) }
942    
943    fn min_at_bt(&self, a: &BackendArray, b: &BackendArray, c: &BackendArray, n: usize, m: usize) -> Result<()>
944    { self.check_and_enqueue_nd_range_for_op("min_at_bt", a, b, c, n, m, 2, 2, true) }
945
946    fn min_a_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
947    { self.check_and_enqueue_nd_range_for_scalar("min_a_b_for_scalar", a, b, c, n, m, 2, 2, true) }
948
949    fn min_at_b_for_scalar(&self, a: &BackendArray, b: f32, c: &BackendArray, n: usize, m: usize) -> Result<()>
950    { self.check_and_enqueue_nd_range_for_scalar("min_at_b_for_scalar", a, b, c, n, m, 2, 2, true) }
951}
952
953#[cfg(test)]
954mod tests;
unmtx_gpu/opencl.rs

unmtx_gpu/
opencl.rs
