ae2f_docs
Slp.auto.h
Go to the documentation of this file.
1#undef __ae2f_MACRO_GENERATED
2#define __ae2f_MACRO_GENERATED 1
3#ifndef Slp_h
4#define Slp_h
5
6#define ae2f_NEED_CLASS 0
7
8#include <ae2fVK/clspv_clkeys.h>
9#undef __ae2f_MACRO_GENERATED
10#define __ae2f_MACRO_GENERATED 1
11#include <ae2f/Ann/Slp.core.h>
12#undef __ae2f_MACRO_GENERATED
13#define __ae2f_MACRO_GENERATED 1
14#include <ae2f/Ann/Slp.auto.h>
15#undef __ae2f_MACRO_GENERATED
16#define __ae2f_MACRO_GENERATED 1
17#include "mac.h"
18#undef __ae2f_MACRO_GENERATED
19#define __ae2f_MACRO_GENERATED 1
20
21typedef ae2f_AnnSlpPredictOne_t _clSlpPredict_t;
22
23#define _clSlpPredict(
24 /** tparam */
25
26
27 /** param */
28 /* , _clSlpPredict_t */ v_predict,
29 /* ae2f_float_t* const */ loc,
30 /* constae2f_float_t* const */ p_inp,
31 /* constae2f_float_t* const */ p_weight,
32 /* constae2f_float_t* const */ p_bias,
33 /* const size_t */ iidx,
34 /* const size_t */ isz,
35 /* const size_t */ oidx,
36 /* const size_t */ osz,
37 /* ae2f_AnnAct_t */ ACT \
38)\
39{
40 if((iidx) == 0 && (oidx) < (osz)) {
41 _clSlpPredict_t _v_predict;
42 (_v_predict).m_tmp = 0;
43 for((_v_predict).m_j = (isz); (_v_predict).m_j--;) {
44 (_v_predict).m_tmp
45 += p_inp[(_v_predict).m_j] * p_weight[(_v_predict).m_j + (isz) * (oidx)];
46 }
47
48 (_v_predict).m_tmp += (p_bias)[oidx];
49 ACT(&(_v_predict).m_ret, (_v_predict).m_tmp);
50 (loc)[oidx] = (_v_predict).m_ret;
51 } \
52}
53
54ae2f_structdef(struct, _clSlpPredict_t_Q) {
55 ae2f_float_t m_sum, m_ret;
56};
57
58/** Quick version, not precise. */
59#define _clSlpPredict_Q(
60 /** tparam */
61
62
63 /** param */
64 /* , _clSlpPredict_t_Q */ v_predict,
65 /* __localae2f_float_t* const */ loc,
66 /* constae2f_float_t* const */ p_inp,
67 /* constae2f_float_t* const */ p_weight,
68 /* constae2f_float_t* const */ p_bias,
69 /* const size_t */ iidx,
70 /* const size_t */ isz,
71 /* const size_t */ oidx,
72 /* const size_t */ osz,
73 /* ae2f_AnnAct_t */ ACT \
74)\
75{
76 if((oidx) < (osz) && (iidx) < (isz)) {
77 (v_predict).m_sum = work_group_reduce_add(
78 (p_weight)[(oidx) * (isz) + (iidx)] * (p_inp)[iidx]
79 );
80
81 if((iidx) == 0) {
82 ACT((&(v_predict).m_ret), ((v_predict).m_sum + (p_bias)[oidx]));
83 (loc)[oidx] = (v_predict).m_ret;
84 }
85 } \
86}
87
88#define clSlpPredict CL_Q_CVRT(_clSlpPredict)
89#define clSlpPredict_t CL_Q_CVRT(_clSlpPredict_t)
90
91#endif
92
93#undef __ae2f_MACRO_GENERATED
94
95#define __ae2f_MACRO_GENERATED 0
ae2f_structdef
#define ae2f_structdef(key, name)
Definition Cast.h:110
ae2f_float_t
ae2f_float ae2f_float_t
Definition Float.h:38
_r_isz
#define _r_isz(lidx)
Definition Mlp.cl.c:54
m_weight
#define m_weight
_r_osz
#define _r_osz(lidx)
Definition Mlp.cl.c:55
kTrainAuto
__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
p_weight
#define p_weight
sizeof(ae2f_float_t) * pgsz_sqr * llsz
Definition Mlp.cl.c:37
_r_weight
#define _r_weight(lidx)
Definition Mlp.cl.c:51
ACT
#define ACT(layer_idx, r, x)
Definition Mlp.cl.c:10
pgsz
#define pgsz
Definition Mlp.cl.c:21
ACT_DERIV
#define ACT_DERIV(layer_idx, r, x)
Definition Mlp.cl.c:14
p_outstream
#define p_outstream
sizeof(ae2f_float_t) * lsz * pgsz
Definition Mlp.cl.c:33
r_out
#define r_out
Definition Mlp.cl.c:58
pgsz_sqr
#define pgsz_sqr
Definition Mlp.cl.c:22
kPredictStream
__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
llsz
#define llsz
Definition Mlp.cl.c:28
m_bias
#define m_bias
_r_bias
#define _r_bias(lidx)
Definition Mlp.cl.c:52
l_inp
#define l_inp(O_R)
Definition Mlp.cl.c:73
l_delta
#define l_delta
Definition Mlp.cl.c:77
r_inp
#define r_inp
Definition Mlp.cl.c:57
l_out
#define l_out(O_R)
Definition Mlp.cl.c:74
l_delta_then
#define l_delta_then
Definition Mlp.cl.c:78
r_weight_then
#define r_weight_then
Definition Mlp.cl.c:67
_r_out
#define _r_out(lidx)
Definition Mlp.cl.c:50
p_deltastream
#define p_deltastream
sizeof(ae2f_float_t) * pgsz * llsz
Definition Mlp.cl.c:43
p_goal
#define p_goal
sizeof(ae2f_float_t) * pgsz
Definition Mlp.cl.c:46
p_bias
#define p_bias
sizeof(ae2f_float_t) * pgsz * llsz
Definition Mlp.cl.c:40
kPredict
__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
STATIC_ASSERT_LRLSZEL_SZ
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
p_layerszlist
#define p_layerszlist
lsz * sizeof(uint32_t)
Definition Mlp.cl.c:31
_r_inp
#define _r_inp(lidx)
Definition Mlp.cl.c:49
r_delta
#define r_delta
Definition Mlp.cl.c:61
_r_delta
#define _r_delta(lidx)
Definition Mlp.cl.c:53
lp_deltastream
#define lp_deltastream
Definition Mlp.cl.c:76
r_weight
#define r_weight
Definition Mlp.cl.c:59
r_isz
#define r_isz
Definition Mlp.cl.c:62
lsz
#define lsz
Definition Mlp.cl.c:25
STATIC_ASSERT_LRLSZ_SZ
char STATIC_ASSERT_LRLSZ_SZ[sizeof(lrlsz_t)==sizeof(lrlszel_t) *3 ? 1 :-1]
Definition Mlp.cl.c:157
r_osz
#define r_osz
Definition Mlp.cl.c:63
r_bias
#define r_bias
Definition Mlp.cl.c:60
kFollow
__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
__global
#define __global
Definition clspv_clkeys.h:22
uint32_t
#define uint32_t
Definition clspv_clkeys.h:37
__local
#define __local
Definition clspv_clkeys.h:24
size_t
#define size_t
Definition clspv_clkeys.h:44
CLK_ALL_MEM_FENCE
#define CLK_ALL_MEM_FENCE
Contains both LOCAL and GLOBAL.
Definition clspv_clkeys.h:69
__kernel
#define __kernel
Definition clspv_clkeys.h:21
get_global_id
size_t get_global_id(uint dimindx)
work_group_reduce_add
#define work_group_reduce_add(x)
Definition clspv_clkeys.h:71
CLK_LOCAL_MEM_FENCE
@ CLK_LOCAL_MEM_FENCE
Definition clspv_clkeys.h:60
get_global_size
size_t get_global_size(uint dimindx)
barrier
void barrier(cl_mem_fence_flags flags)
__ae2f_AnnSlpFollowOneW_imp
#define __ae2f_AnnSlpFollowOneW_imp(inp, delta, weight, learningrate, inp_sz, inp_idx, out_sz, out_idx)
Weights.
Definition Slp.auto.h:366
__ae2f_AnnSlpFollowOneB_imp
#define __ae2f_AnnSlpFollowOneB_imp(r_bias, delta, learningrate_bias)
Definition Slp.auto.h:385
__ae2f_AnnSlpFetchDeltaOne_imp
#define __ae2f_AnnSlpFetchDeltaOne_imp(v_fetchdelta_0, v_fetchdelta_1, out, out_desired, actderiv_opt, lossderiv, retdelta, oidx, osz)
Definition Slp.auto.h:513
CL_Q_CVRT
#define CL_Q_CVRT(x)
Definition mac.h:12
CAST
#define CAST(t, x)
Definition mac.h:5
_clMlpGetHD1
#define _clMlpGetHD1(r_delta_then, i_weight, i_delta, i_iidx, i_isz, i_oidx, i_osz)
Definition Mlp.auto.h:73
_clMlpRvrse
#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
clMlpGetHD
#define clMlpGetHD(...)
GetHidDelta Need no structure.
Definition Mlp.auto.h:120
clMlpGetHD1_t
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
_clMlpGetHD
#define _clMlpGetHD(ONE, r_delta_then, i_weight, i_delta, i_iidx, i_isz, i_oidx, i_osz)
Definition Mlp.auto.h:99
clSlpPredict
#define clSlpPredict
Definition Slp.auto.h:88
_clSlpPredict
#define _clSlpPredict(v_predict, loc, p_inp, p_weight, p_bias, iidx, isz, oidx, osz, ACT)
Definition Slp.auto.h:23
_clSlpPredict_t
ae2f_AnnSlpPredictOne_t _clSlpPredict_t
Definition Slp.auto.h:21
clSlpPredict_t
#define clSlpPredict_t
Definition Slp.auto.h:89
Generated by doxygen 1.14.0