ocl-core 0.11.5

A low-level OpenCL API.
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

static kernel: &'static str = r#"
    __kernel void pyrs_ltp(
                __global uchar const* const axn_states,
                __global uchar const* const cel_states,
                __global uchar const* const cel_tft_best_den_ids,
                __global uchar const* const cel_tft_best_den_states,
                __global uchar const* const den_states,
                __global uchar const* const syn_states,
                __private uint const tfts_per_cel,
                __private uint const dens_per_tft_l2,
                __private uint const syns_per_den_l2,
                __private uint const cels_per_cel_grp,
                __private uint const axn_idz_cel_lyr,
                __private int const learning_rate_l2i,
                __private int const rnd,
                __global uchar* const syn_flag_sets,
                __global uchar* const cel_flag_sets,
                __global int* const aux_ints_0,
                __global int* const aux_ints_1,
                __global char* const syn_strengths)
    {
        uint const cel_grp_id = get_global_id(0);
        uint const cel_grp_count = get_global_size(0);
        uint const cel_count = mul24(cel_grp_count, cels_per_cel_grp);
        uint const cel_idz_cel_grp = mul24(cel_grp_id, cels_per_cel_grp);

         // TODO: (EVALUATE) Make 'cels_per_cel_grp' and 'tfts_per_cel' a constant and unroll loops.
         //    - Will mean making a separate program for each layer of pyramidals.
         //    - Could do more harm than good due to program size bloat.
         //    - Possibly do this for tfts only.
        for (uint cel_id_cel_grp = 0; cel_id_cel_grp < cels_per_cel_grp; cel_id_cel_grp++) {
            uint const cel_idx = cel_idz_cel_grp + cel_id_cel_grp;
            uint const cel_axn_idx = axn_idz_cel_lyr + cel_idx;

            uchar cel_flag_set = cel_flag_sets[cel_idx];

            int const cel_is_concrete = axn_states[cel_axn_idx] != 0;
            int const cel_is_vatic = cel_states[cel_idx] != 0;
            int const cel_prev_concrete = (cel_flag_set & (CEL_PREV_CONCRETE_FLAG)) == (CEL_PREV_CONCRETE_FLAG);
            int const cel_prev_vatic = (cel_flag_set & (CEL_PREV_VATIC_FLAG)) == (CEL_PREV_VATIC_FLAG);
            int const cel_best_in_col = (cel_flag_set & (CEL_BEST_IN_COL_FLAG)) == (CEL_BEST_IN_COL_FLAG);

            for (uint tft_id = 0; tft_id < tfts_per_cel; tft_id++) {
                // uint const cel_tft_idx = mad24(cel_idx, tfts_per_cel, tft_id);
                uint const cel_tft_idx = calc_cel_tft_idx(cel_count, cel_idx, tfts_per_cel, tft_id);
                uint const den_idz_tft = cel_tft_idx << dens_per_tft_l2;

                uchar const den_id_tft_best = cel_tft_best_den_ids[cel_tft_idx];

                uint const syn_idz_tft = den_idz_tft << syns_per_den_l2;
                uint const syn_idz_best_den_tft = (den_idz_tft + den_id_tft_best) << syns_per_den_l2;

                int const tuft_is_active = cel_tft_best_den_states[cel_tft_idx] != 0;

                if (cel_is_concrete) {
                    if (tuft_is_active) {
                        // PREVIOUS (CORRECT) PREDICTION (EVERY PYR IN COL): REINFORCE DEN
                        // ANOMALY (NO PREVIOUS PREDICTION, BEST PYR IN COLUMN ONLY): TRAIN NEW DEN
                        if (cel_prev_vatic | cel_best_in_col) {
                            dst_syns__active__stpot_stdep(syn_states, syn_idz_best_den_tft, syns_per_den_l2, rnd,
                                syn_flag_sets, syn_strengths);
                        }
                    }

                    // TODO: Could be moved into above if block
                    cel_flag_set |= CEL_PREV_CONCRETE_FLAG;

                } else if (cel_prev_concrete) {
                    tft_syns_trm(syn_states, syn_idz_tft, syns_per_den_l2 + dens_per_tft_l2, rnd,
                        learning_rate_l2i, syn_flag_sets, aux_ints_0, syn_strengths);

                    cel_flag_set &= ~CEL_PREV_CONCRETE_FLAG;
                }
            }

            cel_flag_set &= ~CEL_PREV_VATIC_FLAG;
            cel_flag_set |= mul24(cel_is_vatic, CEL_PREV_VATIC_FLAG);

            cel_flag_sets[cel_idx] = cel_flag_set;
        }
    }
"#;


#[test]
fn complex_kernel() {
    unimplemented!();
}