Neuron Net
NeuroNet.cl
Go to the documentation of this file.
1 //---
7 //--- by default some GPU doesn't support doubles
8 //--- cl_khr_fp64 directive is used to enable work with doubles
9 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
10 //+------------------------------------------------------------------+
14 //+------------------------------------------------------------------+
15 __kernel void FeedForward(__global double *matrix_w,
16  __global double *matrix_i,
17  __global double *matrix_o,
18  int inputs,
19  int activation
20  )
21  {
22  int i=get_global_id(0);
23  double sum=0.0;
24  double4 inp, weight;
25  int shift=(inputs+1)*i;
26  for(int k=0; k<=inputs; k=k+4)
27  {
28  switch(inputs-k)
29  {
30  case 0:
31  inp=(double4)(1,0,0,0);
32  weight=(double4)(matrix_w[shift+k],0,0,0);
33  break;
34  case 1:
35  inp=(double4)(matrix_i[k],1,0,0);
36  weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+1],0,0);
37  break;
38  case 2:
39  inp=(double4)(matrix_i[k],matrix_i[k+1],1,0);
40  weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+1],matrix_w[shift+k+2],0);
41  break;
42  case 3:
43  inp=(double4)(matrix_i[k],matrix_i[k+1],matrix_i[k+2],1);
44  weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+1],matrix_w[shift+k+2],matrix_w[shift+k+3]);
45  break;
46  default:
47  inp=(double4)(matrix_i[k],matrix_i[k+1],matrix_i[k+2],matrix_i[k+3]);
48  weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+1],matrix_w[shift+k+2],matrix_w[shift+k+3]);
49  break;
50  }
51  sum+=dot(inp,weight);
52  }
53  switch(activation)
54  {
55  case 0:
56  sum=tanh(sum);
57  break;
58  case 1:
59  sum=1/(1+exp(-clamp(sum,-50.0,50.0)));
60  break;
61  case 2:
62  if(sum<0)
63  sum*=0.01;
64  break;
65  default:
66  break;
67  }
68  matrix_o[i]=sum;
69  }
70 //+------------------------------------------------------------------+
74 //+------------------------------------------------------------------+
75 __kernel void CalcOutputGradient(__global double *matrix_t,
76  __global double *matrix_o,
77  __global double *matrix_ig,
78  int activation
79  )
80  {
81  int i=get_global_id(0);
82  double temp=0;
83  double out=matrix_o[i];
84  switch(activation)
85  {
86  case 0:
87  temp=clamp(matrix_t[i],-1.0,1.0)-out;
88  temp=temp*(1-pow(out==1 || out==-1 ? 0.99999999 : out,2));
89  break;
90  case 1:
91  temp=clamp(matrix_t[i],0.0,1.0)-out;
92  temp=temp*(out==0 || out==1 ? 0.00000001 : (out*(1-out)));
93  break;
94  case 2:
95  temp=(matrix_t[i]-out)*(out>=0 ? 1.0 : 0.01);
96  break;
97  default:
98  temp=(matrix_t[i]-out);
99  break;
100  }
101  matrix_ig[i]=temp;
102  }
103 //+------------------------------------------------------------------+
107 //+------------------------------------------------------------------+
108 __kernel void CalcHiddenGradient(__global double *matrix_w,
109  __global double *matrix_g,
110  __global double *matrix_o,
111  __global double *matrix_ig,
112  int outputs,
113  int activation
114  )
115  {
116  int i=get_global_id(0);
117  int inputs=get_global_size(0);
118  double sum=0;
119  double out=matrix_o[i];
120  double4 grad, weight;
121  for(int k=0;k<outputs;k+=4)
122  {
123  switch(outputs-k)
124  {
125  case 1:
126  grad=(double4)(matrix_g[k],0,0,0);
127  weight=(double4)(matrix_w[k*inputs+i],0,0,0);
128  break;
129  case 2:
130  grad=(double4)(matrix_g[k],matrix_g[k+1],0,0);
131  weight=(double4)(matrix_w[k*inputs+i],matrix_w[k*inputs+i+1],0,0);
132  break;
133  case 3:
134  grad=(double4)(matrix_g[k],matrix_g[k+1],matrix_g[k+2],0);
135  weight=(double4)(matrix_w[k*inputs+i],matrix_w[k*inputs+i+1],matrix_w[k*inputs+i+2],0);
136  break;
137  default:
138  grad=(double4)(matrix_g[k],matrix_g[k+1],matrix_g[k+2],matrix_g[k+3]);
139  weight=(double4)(matrix_w[k*inputs+i],matrix_w[k*inputs+i+1],matrix_w[k*inputs+i+2],matrix_w[k*inputs+i+3]);
140  break;
141  }
142  sum+=dot(grad,weight);
143  }
144  switch(activation)
145  {
146  case 0:
147  sum=clamp(sum+out,-1.0,1.0)-out;
148  sum=sum*(1-pow(out==1 || out==-1 ? 0.99999999 : out,2));
149  break;
150  case 1:
151  sum=clamp(sum+out,0.0,1.0)-out;
152  sum=sum*(out==0 || out==1 ? 0.00000001 : (out*(1-out)));
153  break;
154  case 2:
155  if(out<0)
156  sum*=0.01;
157  break;
158  default:
159  break;
160  }
161  matrix_ig[i]=sum;
162  }
163 //+------------------------------------------------------------------+
167 //+------------------------------------------------------------------+
168 __kernel void UpdateWeightsMomentum(__global double *matrix_w,
169  __global double *matrix_g,
170  __global double *matrix_i,
171  __global double *matrix_dw,
172  int inputs,
173  double learning_rates,
174  double momentum
175  )
176  {
177  int i=get_global_id(0);
178  int j=get_global_id(1);
179  int wi=i*(inputs+1)+j;
180  double delta=learning_rates*matrix_g[i]*(j<inputs ? matrix_i[j] : 1) + momentum*matrix_dw[wi];
181  matrix_dw[wi]=delta;
182  matrix_w[wi]+=delta;
183  };
184 //+------------------------------------------------------------------+
188 //+------------------------------------------------------------------+
189 __kernel void UpdateWeightsAdam(__global double *matrix_w,
190  __global const double *matrix_g,
191  __global const double *matrix_i,
192  __global double *matrix_m,
193  __global double *matrix_v,
194  const int inputs,
195  const double l,
196  const double b1,
197  const double b2
198  )
199  {
200  const int i=get_global_id(0);
201  const int j=get_global_id(1);
202  const int wi=i*(inputs+1)+j*4;
203  double4 m, v, weight, inp;
204  switch(inputs+1-j*4)
205  {
206  case 0:
207  inp=(double4)(1,0,0,0);
208  weight=(double4)(matrix_w[wi],0,0,0);
209  m=(double4)(matrix_m[wi],0,0,0);
210  v=(double4)(matrix_v[wi],0,0,0);
211  break;
212  case 1:
213  inp=(double4)(matrix_i[j],1,0,0);
214  weight=(double4)(matrix_w[wi],matrix_w[wi+1],0,0);
215  m=(double4)(matrix_m[wi],matrix_m[wi+1],0,0);
216  v=(double4)(matrix_v[wi],matrix_v[wi+1],0,0);
217  break;
218  case 2:
219  inp=(double4)(matrix_i[j],matrix_i[j+1],1,0);
220  weight=(double4)(matrix_w[wi],matrix_w[wi+1],matrix_w[wi+2],0);
221  m=(double4)(matrix_m[wi],matrix_m[wi+1],matrix_m[wi+2],0);
222  v=(double4)(matrix_v[wi],matrix_v[wi+1],matrix_v[wi+2],0);
223  break;
224  case 3:
225  inp=(double4)(matrix_i[j],matrix_i[j+1],matrix_i[j+2],1);
226  weight=(double4)(matrix_w[wi],matrix_w[wi+1],matrix_w[wi+2],matrix_w[wi+3]);
227  m=(double4)(matrix_m[wi],matrix_m[wi+1],matrix_m[wi+2],matrix_m[wi+3]);
228  v=(double4)(matrix_v[wi],matrix_v[wi+1],matrix_v[wi+2],matrix_v[wi+3]);
229  break;
230  default:
231  inp=(double4)(matrix_i[j],matrix_i[j+1],matrix_i[j+2],matrix_i[j+3]);
232  weight=(double4)(matrix_w[wi],matrix_w[wi+1],matrix_w[wi+2],matrix_w[wi+3]);
233  m=(double4)(matrix_m[wi],matrix_m[wi+1],matrix_m[wi+2],matrix_m[wi+3]);
234  v=(double4)(matrix_v[wi],matrix_v[wi+1],matrix_v[wi+2],matrix_v[wi+3]);
235  break;
236  }
237  double4 g=(double4)(matrix_g[i])*inp;
238  double4 mt=b1*m+(1-b1)*g;
239  double4 vt=b2*v+(1-b2)*pow(g,2);
240  double4 delta=l*mt/(vt>0 ? sqrt(vt) : l*10);
241  switch(inputs+1-j*4)
242  {
243  case 2:
244  matrix_w[wi+2]+=delta.s2;
245  matrix_m[wi+2]=mt.s2;
246  matrix_v[wi+2]=vt.s2;
247  case 1:
248  matrix_w[wi+1]+=delta.s1;
249  matrix_m[wi+1]=mt.s1;
250  matrix_v[wi+1]=vt.s1;
251  case 0:
252  matrix_w[wi]+=delta.s0;
253  matrix_m[wi]=mt.s0;
254  matrix_v[wi]=vt.s0;
255  break;
256  default:
257  matrix_w[wi]+=delta.s0;
258  matrix_m[wi]=mt.s0;
259  matrix_v[wi]=vt.s0;
260  matrix_w[wi+1]+=delta.s1;
261  matrix_m[wi+1]=mt.s1;
262  matrix_v[wi+1]=vt.s1;
263  matrix_w[wi+2]+=delta.s2;
264  matrix_m[wi+2]=mt.s2;
265  matrix_v[wi+2]=vt.s2;
266  matrix_w[wi+3]+=delta.s3;
267  matrix_m[wi+3]=mt.s3;
268  matrix_v[wi+3]=vt.s3;
269  break;
270  }
271  };
272 //+------------------------------------------------------------------+
275 //+------------------------------------------------------------------+
276 __kernel void FeedForwardProof(__global double *matrix_i,
277  __global double *matrix_o,
278  int inputs,
279  int window,
280  int step
281  )
282  {
283  int i=get_global_id(0);
284  int pos=i*step;
285  double result=matrix_o[pos];
286  for(int k=1; k<window; k=k+1)
287  {
288  int shift=k+pos;
289  if(shift>=inputs)
290  break;
291  result=max(result,matrix_o[shift]);
292  }
293  matrix_o[i]=result;
294  }
295 //+------------------------------------------------------------------+
298 //+------------------------------------------------------------------+
299 __kernel void CalcInputGradientProof(__global double *matrix_i,
300  __global double *matrix_g,
301  __global double *matrix_o,
302  __global double *matrix_ig,
303  int outputs,
304  int window,
305  int step
306  )
307  {
308  int i=get_global_id(0);
309  double prev_gradient=0;
310  double value=matrix_i[i];
311  int start=i-window+step;
312  start=(start-start%step)/step;
313  int stop=(i-i%step)/step+1;
314  for(int out=max(0,start); out<min(outputs,stop); out++)
315  {
316  if(value==matrix_o[out])
317  prev_gradient+=matrix_g[out];
318  }
319  matrix_ig[i]=prev_gradient;
320  }
321 //+------------------------------------------------------------------+
324 //+------------------------------------------------------------------+
325 __kernel void FeedForwardConv(__global double *matrix_w,
326  __global double *matrix_i,
327  __global double *matrix_o,
328  int inputs,
329  int step,
330  int window_in,
331  int window_out,
332  uint activation
333  )
334  {
335  int i=get_global_id(0);
336  int w_in=window_in;
337  int w_out=window_out;
338  double sum=0.0;
339  double4 inp, weight;
340  int shift_out=w_out*i;
341  int shift_in=step*i;
342  for(int out=0;out<w_out;out++)
343  {
344  int shift=(w_in+1)*out;
345  int stop=(w_in<=(inputs-shift_in) ? w_in : (inputs-shift_in));
346  for(int k=0; k<=stop; k=k+4)
347  {
348  switch(stop-k)
349  {
350  case 0:
351  inp=(double4)(1,0,0,0);
352  weight=(double4)(matrix_w[shift+k],0,0,0);
353  break;
354  case 1:
355  inp=(double4)(matrix_i[shift_in+k],1,0,0);
356  weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+1],0,0);
357  break;
358  case 2:
359  inp=(double4)(matrix_i[shift_in+k],matrix_i[shift_in+k+1],1,0);
360  weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+1],matrix_w[shift+k+2],0);
361  break;
362  case 3:
363  inp=(double4)(matrix_i[shift_in+k],matrix_i[shift_in+k+1],matrix_i[shift_in+k+2],1);
364  weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+1],matrix_w[shift+k+2],matrix_w[shift+k+3]);
365  break;
366  default:
367  inp=(double4)(matrix_i[shift_in+k],matrix_i[shift_in+k+1],matrix_i[shift_in+k+2],matrix_i[shift_in+k+3]);
368  weight=(double4)(matrix_w[shift+k],matrix_w[shift+k+1],matrix_w[shift+k+2],matrix_w[shift+k+3]);
369  break;
370  }
371  sum+=dot(inp,weight);
372  }
373  switch(activation)
374  {
375  case 0:
376  sum=tanh(sum);
377  break;
378  case 1:
379  sum=1/(1+exp(-clamp(sum,-50.0,50.0)));
380  break;
381  case 2:
382  if(sum<0)
383  sum*=0.01;
384  break;
385  default:
386  break;
387  }
388  matrix_o[out+shift_out]=sum;
389  }
390  }
391 //+------------------------------------------------------------------+
394 //+------------------------------------------------------------------+
395 __kernel void CalcHiddenGradientConv(__global double *matrix_w,
396  __global double *matrix_g,
397  __global double *matrix_o,
398  __global double *matrix_ig,
399  int outputs,
400  int step,
401  int window_in,
402  int window_out,
403  uint activation
404  )
405  {
406  int i=get_global_id(0);
407  int inputs=get_global_size(0);
408  double sum=0;
409  double out=matrix_o[i];
410  int start=i-window_in+step;
411  start=(start-start%step)/step;
412  int stop=(i-i%step)/step+1;
413  if(stop>(outputs/window_out))
414  stop=outputs;
415  for(int h=0;h<window_out;h+=4)
416  {
417  for(int k=start;k<stop;k++)
418  {
419  int shift_w=(stop-k-1)*step+i%step+h*(window_in+1);
420  int shift_g=k*window_out+h;
421  if(shift_g>=outputs || shift_w>=(window_in+1)*window_out)
422  break;
423  sum+=matrix_g[k*window_out+h]*matrix_w[shift_w];
424  }
425  }
426  switch(activation)
427  {
428  case 0:
429  sum=clamp(sum+out,-1.0,1.0)-out;
430  sum=sum*(1-pow(out==1 || out==-1 ? 0.99999999 : out,2));
431  break;
432  case 1:
433  sum=clamp(sum+out,0.0,1.0)-out;
434  sum=sum*(out==0 || out==1 ? 0.00000001 : (out*(1-out)));
435  break;
436  case 2:
437  if(out<0)
438  sum*=0.01;
439  break;
440  default:
441  break;
442  }
443  matrix_ig[i]=sum;
444  }
445 //+------------------------------------------------------------------+
448 //+------------------------------------------------------------------+
449 __kernel void UpdateWeightsConvMomentum(__global double *matrix_w,
450  __global double *matrix_g,
451  __global double *matrix_i,
452  __global double *matrix_dw,
453  int inputs,
454  double learning_rates,
455  double momentum,
456  int window_in,
457  int window_out,
458  int step
459  )
460  {
461  const int i=get_global_id(0);
462  const int shift=i%(window_in+1);
463  const int shift_out=(i-shift)/(window_in+1);
464  int total=(inputs-window_in)%step;
465  total=(inputs-window_in-total)/step+(total>0 ? 1 : 0);
466  double grad=0;
467  for(int t=0;t<total;t++)
468  {
469  if(shift!=window_in && (shift+t*window_in)>=inputs)
470  break;
471  grad+=matrix_g[t*window_out+shift_out]*(shift==window_in ? 1 : matrix_i[shift+t*step]);
472  }
473  double delta=learning_rates*grad + momentum*matrix_dw[i];
474  matrix_dw[i]=delta;
475  matrix_w[i]+=delta;
476  };
477 //+------------------------------------------------------------------+
480 //+------------------------------------------------------------------+
481 __kernel void UpdateWeightsConvAdam(__global double *matrix_w,
482  __global const double *matrix_g,
483  __global const double *matrix_i,
484  __global double *matrix_m,
485  __global double *matrix_v,
486  const int inputs,
487  const double l,
488  const double b1,
489  const double b2,
490  int window_in,
491  int window_out,
492  int step
493  )
494  {
495  const int i=get_global_id(0);
496  const int shift=i%(window_in+1);
497  const int shift_out=(i-shift)/(window_in+1);
498  int total=(inputs-(window_in-step))%step;
499  total=(inputs-(window_in-step)-total)/step+(total>0 ? 1 : 0);
500  double grad=0;
501  for(int t=0;t<total;t++)
502  {
503  if(shift!=window_in && (shift+t*window_in)>=inputs)
504  break;
505  grad+=matrix_g[t*window_out+shift_out]*(shift==window_in ? 1 : matrix_i[shift+t*step]);
506  }
507  double mt=b1*matrix_m[i]+(1-b1)*grad;
508  double vt=b2*matrix_v[i]+(1-b2)*pow(grad,2);
509  double delta=(vt>0 ? l*mt/sqrt(vt) : 0);
510  matrix_w[i]+=delta;
511  matrix_m[i]=mt;
512  matrix_v[i]=vt;
513  };
514 //+------------------------------------------------------------------+
518 //+------------------------------------------------------------------+
519 __kernel void AttentionScore(__global double *querys,
520  __global double *keys,
521  __global double *score,
522  int dimension
523  )
524  {
525  int q=get_global_id(0);
526  int shift_q=q*dimension;
527  int units=get_global_size(0);
528  int shift_s=q*units;
529  double koef=sqrt((double)(units*dimension));
530  if(koef<1)
531  koef=1;
532  double sum=0;
533  for(int k=0;k<units;k++)
534  {
535  double result=0;
536  int shift_k=k*dimension;
537  for(int i=0;i<dimension;i++)
538  result+=(querys[shift_q+i]*keys[shift_k+i]);
539  result=exp(clamp(result/koef,-500.0,500.0));
540  score[shift_s+k]=result;
541  sum+=result;
542  }
543  for(int k=0;k<units;k++)
544  score[shift_s+k]/=sum;
545  }
546 //+------------------------------------------------------------------+
550 //+------------------------------------------------------------------+
551 __kernel void AttentionOut(__global double *scores,
552  __global double *values,
553  __global double *inputs,
554  __global double *out
555  )
556  {
557  int units=get_global_size(0);
558  int u=get_global_id(0);
559  int d=get_global_id(1);
560  int dimension=get_global_size(1);
561  int shift=u*dimension+d;
562  double result=0;
563  for(int i=0;i<units;i++)
564  result+=scores[u*units+i]*values[i*dimension+d];
565  out[shift]=result+inputs[shift];
566  }
567 //+------------------------------------------------------------------+
571 //+------------------------------------------------------------------+
572 __kernel void SumMatrix(__global double *matrix1,
573  __global double *matrix2,
574  __global double *matrix_out,
575  int dimension,
576  double multiplyer
577  )
578  {
579  const int i=get_global_id(0)*dimension;
580  for(int k=0;k<dimension;k++)
581  matrix_out[i+k]=(matrix1[i+k]+matrix2[i+k])*multiplyer;
582  }
583 //+------------------------------------------------------------------+
595 //+------------------------------------------------------------------+
596 __kernel void AttentionIsideGradients(__global double *querys,__global double *querys_g,
597  __global double *keys,__global double *keys_g,
598  __global double *values,__global double *values_g,
599  __global double *scores,
600  __global double *gradient)
601  {
602  int u=get_global_id(0);
603  int d=get_global_id(1);
604  int units=get_global_size(0);
605  int dimension=get_global_size(1);
606  double koef=sqrt((double)(units*dimension));
607  if(koef<1)
608  koef=1;
609 
610  double vg=0;
611  double qg=0;
612  double kg=0;
613  for(int iu=0;iu<units;iu++)
614  {
615  double g=gradient[iu*dimension+d];
616  double sc=scores[iu*units+u];
617  vg+=sc*g;
618 //---
619  double sqg=0;
620  double skg=0;
621  for(int id=0;id<dimension;id++)
622  {
623  sqg+=values[iu*dimension+id]*gradient[u*dimension+id];
624  skg+=values[u*dimension+id]*gradient[iu*dimension+id];
625  }
626  qg+=(scores[u*units+iu]==0 || scores[u*units+iu]==1 ? 0.0001 : scores[u*units+iu]*(1-scores[u*units+iu]))*sqg*keys[iu*dimension+d]/koef;
627 //---
628  kg+=(scores[iu*units+u]==0 || scores[iu*units+u]==1 ? 0.0001 : scores[iu*units+u]*(1-scores[iu*units+u]))*skg*querys[iu*dimension+d]/koef;
629  }
630  int shift=u*dimension+d;
631  values_g[shift]=vg;
632  querys_g[shift]=qg;
633  keys_g[shift]=kg;
634  }
635 //+------------------------------------------------------------------+
641 //+------------------------------------------------------------------+
642 __kernel void Normalize(__global double *buffer,
643  int dimension)
644  {
645  int n=get_global_id(0);
646  int shift=n*dimension;
647  double mean=0;
648  for(int i=0;i<dimension;i++)
649  mean+=buffer[shift+i];
650  mean/=dimension;
651  double variance=0;
652  for(int i=0;i<dimension;i++)
653  variance+=pow(buffer[shift+i]-mean,2);
654  variance=sqrt(variance/dimension);
655  for(int i=0;i<dimension;i++)
656  buffer[shift+i]=(buffer[shift+i]-mean)/(variance==0 ? 1 : variance);
657  }
658 //+------------------------------------------------------------------+
664 //+------------------------------------------------------------------+
665 __kernel void NormalizeWeights(__global double *buffer,
666  int dimension)
667  {
668  int n=get_global_id(0);
669  int shift=n*dimension;
670  double sum=0;
671  for(int i=0;i<dimension;i++)
672  sum=pow(buffer[shift+i],2);
673  sum=sqrt(sum);
674  for(int i=0;i<dimension;i++)
675  buffer[shift+i]/=sum;
676  }
677 //+------------------------------------------------------------------+
b2
#define b2
Second momentum multiplier of Adam optimization.
Definition: NeuroNet.mqh:38
CalcHiddenGradient
__kernel void CalcHiddenGradient(__global double *matrix_w, __global double *matrix_g, __global double *matrix_o, __global double *matrix_ig, int outputs, int activation)
Describes the process of hidden gradients calculation for the Neuron Base (CNeuronBaseOCL).
Definition: NeuroNet.cl:108
UpdateWeightsConvMomentum
__kernel void UpdateWeightsConvMomentum(__global double *matrix_w, __global double *matrix_g, __global double *matrix_i, __global double *matrix_dw, int inputs, double learning_rates, double momentum, int window_in, int window_out, int step)
Describes the process of SGD optimization weights for the Convolution Neuron (CNeuronConvOCL).
Definition: NeuroNet.cl:449
CalcOutputGradient
__kernel void CalcOutputGradient(__global double *matrix_t, __global double *matrix_o, __global double *matrix_ig, int activation)
Describes the process of output gradients calculation for the Neuron Base (CNeuronBaseOCL).
Definition: NeuroNet.cl:75
FeedForward
__kernel void FeedForward(__global double *matrix_w, __global double *matrix_i, __global double *matrix_o, int inputs, int activation)
Describes the forward path process for the Neuron Base (CNeuronBaseOCL).
Definition: NeuroNet.cl:15
NormalizeWeights
__kernel void NormalizeWeights(__global double *buffer, int dimension)
Describes the process of weights matrix normalization.
Definition: NeuroNet.cl:665
UpdateWeightsMomentum
__kernel void UpdateWeightsMomentum(__global double *matrix_w, __global double *matrix_g, __global double *matrix_i, __global double *matrix_dw, int inputs, double learning_rates, double momentum)
Describes the process of SGD optimization weights for the Neuron Base (CNeuronBaseOCL).
Definition: NeuroNet.cl:168
b1
#define b1
First momentum multiplier of Adam optimization.
Definition: NeuroNet.mqh:36
CalcHiddenGradientConv
__kernel void CalcHiddenGradientConv(__global double *matrix_w, __global double *matrix_g, __global double *matrix_o, __global double *matrix_ig, int outputs, int step, int window_in, int window_out, uint activation)
Kernel of the Convolution neuron to transfer gradient to previous layer (CNeuronConvOCL)
Definition: NeuroNet.cl:395
AttentionIsideGradients
__kernel void AttentionIsideGradients(__global double *querys, __global double *querys_g, __global double *keys, __global double *keys_g, __global double *values, __global double *values_g, __global double *scores, __global double *gradient)
' Describes the gradients calculation process for the Neuron of attention layer (CNeuronAttentionOCL)...
Definition: NeuroNet.cl:596
AttentionScore
__kernel void AttentionScore(__global double *querys, __global double *keys, __global double *score, int dimension)
| Describes the Score calculation process for the Neuron of attention layer (CNeuronAttentionOCL).
Definition: NeuroNet.cl:519
UpdateWeightsAdam
__kernel void UpdateWeightsAdam(__global double *matrix_w, __global const double *matrix_g, __global const double *matrix_i, __global double *matrix_m, __global double *matrix_v, const int inputs, const double l, const double b1, const double b2)
Describes the process of Adam optimization weights for the Neuron Base (CNeuronBaseOCL).
Definition: NeuroNet.cl:189
FeedForwardConv
__kernel void FeedForwardConv(__global double *matrix_w, __global double *matrix_i, __global double *matrix_o, int inputs, int step, int window_in, int window_out, uint activation)
Kernel of the Convolution neuron for Feed forward process (CNeuronConvOCL)
Definition: NeuroNet.cl:325
momentum
#define momentum
momentum for SGD optimization
Definition: NeuroNet.mqh:33
UpdateWeightsConvAdam
__kernel void UpdateWeightsConvAdam(__global double *matrix_w, __global const double *matrix_g, __global const double *matrix_i, __global double *matrix_m, __global double *matrix_v, const int inputs, const double l, const double b1, const double b2, int window_in, int window_out, int step)
Describes the process of Adam optimization weights for the Convolution Neuron (CNeuronConvOCL).
Definition: NeuroNet.cl:481
SumMatrix
__kernel void SumMatrix(__global double *matrix1, __global double *matrix2, __global double *matrix_out, int dimension, double multiplyer)
2 . Describes the calculation Sum of 2 matrixs.
Definition: NeuroNet.cl:572
Normalize
__kernel void Normalize(__global double *buffer, int dimension)
Describes the process of matrix normalization.
Definition: NeuroNet.cl:642
CalcInputGradientProof
__kernel void CalcInputGradientProof(__global double *matrix_i, __global double *matrix_g, __global double *matrix_o, __global double *matrix_ig, int outputs, int window, int step)
Kernel of the Pooling neuron to transfer gradient to previous layer (CNeuronProofOCL)
Definition: NeuroNet.cl:299
AttentionOut
__kernel void AttentionOut(__global double *scores, __global double *values, __global double *inputs, __global double *out)
Describes the Attention out calculation process for the Neuron of attention layer (CNeuronAttentionOC...
Definition: NeuroNet.cl:551
FeedForwardProof
__kernel void FeedForwardProof(__global double *matrix_i, __global double *matrix_o, int inputs, int window, int step)
Kernel of the Pooling neuron for Feed forward process (CNeuronProofOCL)
Definition: NeuroNet.cl:276