ha-ndarray 0.3.0

A hardware-accelerated n-dimensional array
Documentation 
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
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167

use memoize::memoize;
use ocl::Program;

use crate::Error;

use super::build;

#[memoize]
pub fn cast(i_type: &'static str, o_type: &'static str) -> Result<Program, Error> {
    let src = format!(
        r#"
        __kernel void cast(
            __global const {i_type}* restrict input,
            __global {o_type}* restrict output)
        {{
            const ulong offset = get_global_id(0);
            output[offset] = ({o_type}) input[offset];
        }}
        "#,
    );

    build(&src)
}

#[memoize]
pub fn dual_boolean(c_type: &'static str, op: &'static str) -> Result<Program, Error> {
    let src = format!(
        r#"
        inline uchar eq(const {c_type} left, const {c_type} right) {{
            return left == right;
        }}

        inline uchar ge(const {c_type} left, const {c_type} right) {{
            return left >= right;
        }}

        inline uchar gt(const {c_type} left, const {c_type} right) {{
            return left > right;
        }}

        inline uchar le(const {c_type} left, const {c_type} right) {{
            return left <= right;
        }}

        inline uchar lt(const {c_type} left, const {c_type} right) {{
            return left < right;
        }}

        inline uchar ne(const {c_type} left, const {c_type} right) {{
            return left != right;
        }}

        inline uchar and(const {c_type} left, const {c_type} right) {{
            return (left != 0) && (right != 0);
        }}

        inline uchar or(const {c_type} left, const {c_type} right) {{
            return (left != 0) || (right != 0);
        }}

        inline uchar xor(const {c_type} left, const {c_type} right) {{
            return (left != 0) ^ (right != 0);
        }}

        __kernel void dual(
            __global const {c_type}* restrict left,
            __global const {c_type}* restrict right,
            __global uchar* restrict output)
        {{
            const ulong offset = get_global_id(0);
            output[offset] = {op}(left[offset], right[offset]);
        }}

        __kernel void dual_scalar(
            __global const {c_type}* restrict left,
            const {c_type} right,
            __global uchar* restrict output)
        {{
            const ulong offset = get_global_id(0);
            output[offset] = {op}(left[offset], right);
        }}
        "#
    );

    build(&src)
}

#[memoize]
pub fn dual(c_type: &'static str, op: &'static str) -> Result<Program, Error> {
    let src = format!(
        r#"
        inline {c_type} _log(const double left, const double right) {{
            return log(left) / log(right);
        }}

        inline {c_type} add(const {c_type} left, const {c_type} right) {{
            return left + right;
        }}

        inline {c_type} div(const {c_type} left, const {c_type} right) {{
            if (right == 0) {{
                return 0;
            }} else {{
                return left / right;
            }}
        }}

        inline {c_type} mul(const {c_type} left, const {c_type} right) {{
            return left * right;
        }}

        inline {c_type} sub(const {c_type} left, const {c_type} right) {{
            return left - right;
        }}

        __kernel void dual(
            __global const {c_type}* restrict left,
            __global const {c_type}* restrict right,
            __global {c_type}* restrict output)
        {{
            const ulong offset = get_global_id(0);
            output[offset] = {op}(left[offset], right[offset]);
        }}

        __kernel void dual_scalar(
            __global const {c_type}* restrict left,
            const {c_type} right,
            __global {c_type}* restrict output)
        {{
            const ulong offset = get_global_id(0);
            output[offset] = {op}(left[offset], right);
        }}
        "#,
    );

    build(&src)
}

pub fn unary(
    f_type: &'static str,
    i_type: &'static str,
    o_type: &'static str,
    op: &'static str,
) -> Result<Program, Error> {
    let src = format!(
        r#"
        inline uchar not(const {i_type} input) {{
            if (input == 0) {{
                return 1;
            }} else {{
                return 0;
            }}
        }}

        inline {f_type} _log(const {f_type} input) {{
            return log(input);
        }}

        __kernel void unary(__global const {i_type}* input, __global {o_type}* output) {{
            const ulong offset = get_global_id(0);
            output[offset] = {op}(input[offset]);
        }}
        "#,
    );

    build(&src)
}