File size: 1,375 Bytes
a8baeed | 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 | use super::{GpuContext, Result};
use ocl::Buffer;
pub(super) const KERNEL_SRC: &str = r#"
__kernel void add_u8(__global const uchar* a,
__global const uchar* b,
__global uchar* output,
const uint count) {
size_t i = get_global_id(0);
if (i >= count) {
return;
}
output[i] = (uchar)(a[i] + b[i]);
}
"#;
/// Element-wise wrapping `u8` addition: `output[i] = a[i] + b[i]`.
/// On Booleans this is a counting sum (0, 1 or 2), so `Negate(Add(x, y))` is NOR.
#[derive(Clone, Copy, Debug, Default)]
pub struct Add;
impl Add {
pub fn launch(
&self,
ctx: &GpuContext,
a: &Buffer<u8>,
b: &Buffer<u8>,
output: &Buffer<u8>,
count: usize,
) -> Result<()> {
let (global, local) = ctx.work_sizes(
count,
&[("a", a.len()), ("b", b.len()), ("output", output.len())],
)?;
if count == 0 {
return Ok(());
}
let kernel = ctx
.proque()
.kernel_builder("add_u8")
.global_work_size(global)
.local_work_size(local)
.arg(a)
.arg(b)
.arg(output)
.arg(count as u32)
.build()?;
// SAFETY: see Negate::launch.
unsafe { kernel.enq()? };
Ok(())
}
}
|