ae2f_docs
Mlp.auto.h
Go to the documentation of this file.
1#undef __ae2f_MACRO_GENERATED
2#define __ae2f_MACRO_GENERATED 1
3#ifndef Mlp_h
4#define Mlp_h
5
6#include "./Slp.auto.h"
7#undef __ae2f_MACRO_GENERATED
8#define __ae2f_MACRO_GENERATED 1
9
10typedef void clMlpGetHD1_t(
11 ae2f_float_t* const r_delta,
12
13 const ae2f_float_t* const i_weight,
14 const ae2f_float_t* const i_delta,
15
16 const size_t i_iidx,
17 const size_t i_isz,
18 const size_t i_oidx,
19 const size_t i_osz
20 );
21
22/**
23 * @brief
24 * delta to delta
25 * Propagate
26 *
27 * @inp aka out_then
28 * @deltaseed
29 * */
30#define _clMlpRvrse(
31 /** tparam */
32
33
34 /** param */
35 /* , ae2f_float_t */ v_tmp,
36 /* ae2f_float_t* const */ r_delta_then,
37 /* const size_t */ i_oidx,
38 /* const size_t */ i_iidx,
39 /* const size_t */ i_isz,
40 /* ae2f_AnnAct_t */ i_actderiv_then,
41 /* constae2f_float_t* const */ i_inp,
42 /* constae2f_float_t* const */ i_deltaseed \
43)\
44{
45 if((i_iidx) < (i_isz) && !(i_oidx)) {
46 i_actderiv_then(&(v_tmp), (i_inp)[i_iidx]);
47 (r_delta_then)[i_iidx] = (v_tmp) * (i_deltaseed)[i_iidx];
48 } \
49}
50
51#define _clMlpGetHD1_Q(
52 /** tparam */
53
54
55 /** param */
56 /* ,ae2f_float_t* const */ r_delta_then,
57 /* constae2f_float_t* const */ i_weight,
58 /* constae2f_float_t* const */ i_delta,
59 /* const size_t */ i_iidx,
60 /* const size_t */ i_isz,
61 /* const size_t */ i_oidx,
62 /* const size_t */ i_osz \
63)\
64{
65 if((i_oidx) < (i_osz) && (i_iidx) < (i_isz)) {
66 (r_delta_then)[(i_iidx)]
68 (i_weight)[(i_isz) * (i_oidx) + (i_iidx)] * (i_delta)[i_oidx]
69 );
70 } \
71}
72
73#define _clMlpGetHD1(
74 /** tparam */
75
76
77 /** param */
78 /* ,ae2f_float_t* const */ r_delta_then,
79 /* constae2f_float_t* const */ i_weight,
80 /* constae2f_float_t* const */ i_delta,
81 /* const size_t */ i_iidx,
82 /* const size_t */ i_isz,
83 /* const size_t */ i_oidx,
84 /* const size_t */ i_osz \
85)\
86{
87 if((i_oidx) == 0 && (i_iidx) < (i_isz)) {
88 size_t v_oidx = (i_osz);
89 ae2f_float_t v_ret = 0;
90
91 while((v_oidx)--) {
92 (v_ret) += (i_weight)[(i_isz) * (v_oidx) + (i_iidx)] * (i_delta)[(v_oidx)];
93 }
94
95 (r_delta_then)[(i_iidx)] = (v_ret);
96 } \
97}
98
99#define _clMlpGetHD(
100 /** tparam */
101
102
103 /** param */
104 /* , clMlpGetHD1_t */ ONE,
105 /* ae2f_float_t* const */ r_delta_then,
106 /* constae2f_float_t* const */ i_weight,
107 /* constae2f_float_t* const */ i_delta,
108 /* const size_t */ i_iidx,
109 /* const size_t */ i_isz,
110 /* const size_t */ i_oidx,
111 /* const size_t */ i_osz \
112)\
113{
114 if((i_iidx) < (i_isz) && (i_oidx) < (i_osz)) {
115 ONE(r_delta_then, i_weight, i_delta, i_iidx, i_isz, i_oidx, i_osz);
116 } \
117}
118
119/** @brief GetHidDelta Need no structure. */
120#define clMlpGetHD(...) _clMlpGetHD(CL_Q_CVRT(_clMlpGetHD1), __VA_ARGS__)
121
122
123#endif
124
125#undef __ae2f_MACRO_GENERATED
126
127#define __ae2f_MACRO_GENERATED 0
#define ae2f_structdef(key, name)
Definition Cast.h:110
ae2f_float ae2f_float_t
Definition Float.h:38
#define _r_isz(lidx)
Definition Mlp.cl.c:54
#define m_weight
#define _r_osz(lidx)
Definition Mlp.cl.c:55
__kernel void kTrainAuto(__global void *glob, __local ae2f_float_t *loc, lrlsz_t lr)
loc ae2f_float_t[lsz - 1][Page]: OutStream ae2f_float_t[lsz - 1][Page]: DeltaStream
Definition Mlp.cl.c:290
#define p_weight
sizeof(ae2f_float_t) * pgsz_sqr * llsz
Definition Mlp.cl.c:37
#define _r_weight(lidx)
Definition Mlp.cl.c:51
#define ACT(layer_idx, r, x)
Definition Mlp.cl.c:10
#define pgsz
Definition Mlp.cl.c:21
#define ACT_DERIV(layer_idx, r, x)
Definition Mlp.cl.c:14
#define p_outstream
sizeof(ae2f_float_t) * lsz * pgsz
Definition Mlp.cl.c:33
#define r_out
Definition Mlp.cl.c:58
#define pgsz_sqr
Definition Mlp.cl.c:22
__kernel void kPredictStream(__global void *glob, __local ae2f_float_t *loc, const uint32_t lsz)
loc ae2f_float_t[Page]: inp ae2f_float_t[Page]: out
Definition Mlp.cl.c:118
#define llsz
Definition Mlp.cl.c:28
#define m_bias
#define _r_bias(lidx)
Definition Mlp.cl.c:52
#define l_inp(O_R)
Definition Mlp.cl.c:73
#define l_delta
Definition Mlp.cl.c:77
#define r_inp
Definition Mlp.cl.c:57
#define l_out(O_R)
Definition Mlp.cl.c:74
#define l_delta_then
Definition Mlp.cl.c:78
#define r_weight_then
Definition Mlp.cl.c:67
#define _r_out(lidx)
Definition Mlp.cl.c:50
#define p_deltastream
sizeof(ae2f_float_t) * pgsz * llsz
Definition Mlp.cl.c:43
#define p_goal
sizeof(ae2f_float_t) * pgsz
Definition Mlp.cl.c:46
#define p_bias
sizeof(ae2f_float_t) * pgsz * llsz
Definition Mlp.cl.c:40
__kernel void kPredict(__global void *glob, __local ae2f_float_t *loc, const uint32_t lsz)
loc ae2f_float_t[Page]: inp ae2f_float_t[Page]: out
Definition Mlp.cl.c:93
char STATIC_ASSERT_LRLSZEL_SZ[sizeof(lrlszel_t)==(sizeof(uint32_t) > sizeof(ae2f_float_t) ? sizeof(uint32_t) :sizeof(ae2f_float_t)) ? 1 :-1]
Definition Mlp.cl.c:155
#define p_layerszlist
lsz * sizeof(uint32_t)
Definition Mlp.cl.c:31
#define _r_inp(lidx)
Definition Mlp.cl.c:49
#define r_delta
Definition Mlp.cl.c:61
#define _r_delta(lidx)
Definition Mlp.cl.c:53
#define lp_deltastream
Definition Mlp.cl.c:76
#define r_weight
Definition Mlp.cl.c:59
#define r_isz
Definition Mlp.cl.c:62
#define lsz
Definition Mlp.cl.c:25
char STATIC_ASSERT_LRLSZ_SZ[sizeof(lrlsz_t)==sizeof(lrlszel_t) *3 ? 1 :-1]
Definition Mlp.cl.c:157
#define r_osz
Definition Mlp.cl.c:63
#define r_bias
Definition Mlp.cl.c:60
__kernel void kFollow(__global void *glob, __local ae2f_float_t *loc, lrlsz_t lr)
loc ae2f_float_t[lsz - 1][Page]: OutStream ae2f_float_t[lsz - 1][Page]: DeltaStream
Definition Mlp.cl.c:166
#define __global
#define uint32_t
#define __local
#define size_t
#define CLK_ALL_MEM_FENCE
Contains both LOCAL and GLOBAL.
#define __kernel
size_t get_global_id(uint dimindx)
#define work_group_reduce_add(x)
@ CLK_LOCAL_MEM_FENCE
size_t get_global_size(uint dimindx)
void barrier(cl_mem_fence_flags flags)
#define __ae2f_AnnSlpFollowOneW_imp(inp, delta, weight, learningrate, inp_sz, inp_idx, out_sz, out_idx)
Weights.
Definition Slp.auto.h:366
#define __ae2f_AnnSlpFollowOneB_imp(r_bias, delta, learningrate_bias)
Definition Slp.auto.h:385
#define __ae2f_AnnSlpFetchDeltaOne_imp(v_fetchdelta_0, v_fetchdelta_1, out, out_desired, actderiv_opt, lossderiv, retdelta, oidx, osz)
Definition Slp.auto.h:513
#define CL_Q_CVRT(x)
Definition mac.h:12
#define CAST(t, x)
Definition mac.h:5
#define _clMlpGetHD1(r_delta_then, i_weight, i_delta, i_iidx, i_isz, i_oidx, i_osz)
Definition Mlp.auto.h:73
#define _clMlpRvrse(v_tmp, r_delta_then, i_oidx, i_iidx, i_isz, i_actderiv_then, i_inp, i_deltaseed)
delta to delta Propagate
Definition Mlp.auto.h:30
#define clMlpGetHD(...)
GetHidDelta Need no structure.
Definition Mlp.auto.h:120
void clMlpGetHD1_t(ae2f_float_t *const r_delta, const ae2f_float_t *const i_weight, const ae2f_float_t *const i_delta, const size_t i_iidx, const size_t i_isz, const size_t i_oidx, const size_t i_osz)
Definition Mlp.auto.h:10
#define _clMlpGetHD(ONE, r_delta_then, i_weight, i_delta, i_iidx, i_isz, i_oidx, i_osz)
Definition Mlp.auto.h:99
#define clSlpPredict
Definition Slp.auto.h:88
#define clSlpPredict_t
Definition Slp.auto.h:89