kelmut/NeuroBook
Watch
1
0
Fork
You've already forked NeuroBook
 0
forked from rosh/NeuroBook
Code Pull requests Activity
NeuroBook/Include/realization/opencl_program.cl

1326 lines
114 KiB
Common Lisp
RawPermalinkNormal ViewHistory

convert
2025-05-30 16:12:30 +02:00
<EFBFBD><EFBFBD>//+------------------------------------------------------------------+
//| opencl_program.cl |
//| Copyright 2021, MetaQuotes Ltd. |
//| https://www.mql5.com |
//+------------------------------------------------------------------+
//--- by default some GPU doesn't support TYPEs
//--- cl_khr_fp64 directive is used to enable work with TYPEs
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
//+------------------------------------------------------------------+
//| Activation functions |
//+------------------------------------------------------------------+
//| Linear activation function |
//| Parameter 'value' Weighted sum of initial data |
//| 'a' defines the angle of inclination of the line |
//| 'b' - the vertical offset of the line |
//+------------------------------------------------------------------+
__kernel void LineActivation(__global TYPE* inputs,
__global TYPE* outputs,
const TYPE a, const TYPE b)
{
size_t i = get_global_id(0);
outputs[i] = (a * inputs[i] + b);
}
//+------------------------------------------------------------------+
//| Sigmoid activation function |
//| Parameter 'value' Weighted sum of initial data |
//| 'a' stretches the range of values of the function |
//| from '0' to 'a' |
//| 'b' shifts the resulting value |
//+------------------------------------------------------------------+
__kernel void SigmoidActivation(__global TYPE* inputs,
__global TYPE* outputs,
const TYPE a, const TYPE b)
{
size_t i = get_global_id(0);
outputs[i] = a / (1 + exp(-inputs[i])) - b;
}
//+------------------------------------------------------------------+
//| Derivative of Sigmoid activation function |
//| Parameter 'value' current point (result of feed forward). |
//| 'a' stretches the range of values of the function from |
//| '0' to 'a' |
//| 'b' shifts the resulting value |
//+------------------------------------------------------------------+
__kernel void SigmoidDerivative(__global TYPE* outputs,
__global TYPE* output_gr,
__global TYPE* input_gr,
const TYPE a, const TYPE b
)
{
size_t i = get_global_id(0);
if(a == 0)
input_gr[i] = 0;
else
{
TYPE z = clamp(outputs[i] + b, (TYPE)0, a);
input_gr[i] = z * (1 - z / a) * output_gr[i];
}
}
//+------------------------------------------------------------------+
//| TANH activation function |
//| Parameter 'value' Weighted sum of initial data |
//+------------------------------------------------------------------+
__kernel void TanhActivation(__global TYPE* inputs,
__global TYPE* outputs)
{
size_t i = get_global_id(0);
outputs[i] = tanh(inputs[i]);
}
//+------------------------------------------------------------------+
//| Derivative of TANH activation function |
//| Parameter 'value' current point (result of feed forward). |
//+------------------------------------------------------------------+
__kernel void TanhDerivative(__global TYPE* outputs,
__global TYPE* output_gr,
__global TYPE* input_gr
)
{
size_t i = get_global_id(0);
input_gr[i] = (1 - pow(outputs[i], 2)) * output_gr[i];
}
//+------------------------------------------------------------------+
//| LReLU activation function |
//| Parameter 'value' current point (result of feed forward). |
//| 'a' leak parameter |
//+------------------------------------------------------------------+
__kernel void LReLUActivation(__global TYPE* inputs,
__global TYPE* outputs,
const TYPE a)
{
size_t i = get_global_id(0);
TYPE value = inputs[i];
outputs[i] = (value > 0 ? value : a * value);
}
//+------------------------------------------------------------------+
//| Derivative of LReLU activation function |
//| Parameter 'value' current point (result of feed forward). |
//| 'a' leak parameter |
//+------------------------------------------------------------------+
__kernel void LReLUDerivative(__global TYPE* outputs,
__global TYPE* output_gr,
__global TYPE* input_gr,
const TYPE a)
{
size_t i = get_global_id(0);
input_gr[i] = (outputs[i] > 0 ? (TYPE)1 : a) * output_gr[i];
}
//+------------------------------------------------------------------+
//| Swish activation function |
//| Parameter 'value' Weighted sum of initial data |
//| 'b' affects the nonlinearity of the function |
//+------------------------------------------------------------------+
__kernel void SwishActivation(__global TYPE* inputs,
__global TYPE* outputs,
const TYPE b)
{
size_t i = get_global_id(0);
TYPE value = inputs[i];
outputs[i] = value / (1 + exp(-b * value));
}
//+------------------------------------------------------------------+
//| Derivative of Swish activation function |
//| Parameter 'value' current point (result of feed forward). |
//| 'value_input' Weighted sum of initial data |
//| 'b' affects the nonlinearity of the function |
//+------------------------------------------------------------------+
__kernel void SwishDerivative(__global TYPE* outputs,
__global TYPE* output_gr,
__global TYPE* input_gr,
const TYPE b,
__global TYPE* inputs)
{
size_t i = get_global_id(0);
TYPE by = b * outputs[i];
input_gr[i] = (by + (1 - by) / (1 + exp(-b * inputs[i]))) * output_gr[i];
}
//+------------------------------------------------------------------+
//| Transfer 4 elements of TYPE vector to TYPE4 |
//| Parameter 'array' source array of data |
//| 'start' first position to copy |
//| 'step' step between elements to copy |
//| 'size' Size of source array |
//| 'shift' Shift in source array to the 1-st copied element |
//+------------------------------------------------------------------+
TYPE4 ToVect4(__global TYPE *array, int start, int step, int size, int shift)
{
TYPE4 result = (TYPE4)(0, 0, 0, 0);
step = max(1, step);
int st = start * step + shift;
if(st < size)
{
int k = (size - shift + step - 1) / step;
switch(k)
{
case 0:
break;
case 1:
result = (TYPE4)(array[st], 0, 0, 0);
break;
case 2:
result = (TYPE4)(array[st], array[st + step], 0, 0);
break;
case 3:
result = (TYPE4)(array[st], array[st + step], array[st + 2 * step], 0);
break;
default:
result = (TYPE4)(array[st], array[st + step], array[st + 2 * step], array[st + 3 * step]);
break;
}
}
return result;
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
TYPE Max4(TYPE4 vect, TYPE value)
{
TYPE result = fmax(vect.s0, value);
result = fmax(vect.s1, result);
result = fmax(vect.s2, result);
return(fmax(vect.s3, result));
}
//+------------------------------------------------------------------+
//| Transfer TYPE4 to 4 elements of TYPE vector |
//| Parameter 'array' target array of data |
//| 'value' source TYPE4 vector |
//| 'start' first position to copy in target array |
//| 'step' step between elements in target array |
//| 'size' Size of target array |
//| 'shift' Shift in target array to the 1-st copied element |
//+------------------------------------------------------------------+
void D4ToArray(__global TYPE *array, TYPE4 value, int start, int step, int size, int shift)
{
step = max(1, step);
int st = start * step + shift;
if(st < size)
{
int k = (size - shift) % step;
k = (size - shift - k) / step - start + (k > 0 ? 1 : 0);
switch(k)
{
case 3:
array[st + 2 * step] = value.s2;
case 2:
array[st + step] = value.s1;
case 1:
array[st] = value.s0;
break;
default:
array[st + 3 * step] = value.s3;
array[st + 2 * step] = value.s2;
array[st + step] = value.s1;
array[st] = value.s0;
break;
}
}
return;
}
//+------------------------------------------------------------------+
//| Kernel Feed Forward of perceptron |
//| Parameter 'inputs' vector of inputs data |
//| 'weights' matrix of weights |
//| 'outputs' output data |
//| 'inputs_total' size of inputs vector |
//+------------------------------------------------------------------+
__kernel void PerceptronFeedForward(__global TYPE *inputs,
__global TYPE *weights,
__global TYPE *outputs,
int inputs_total)
{
const int n = get_global_id(0);
const int weights_total = get_global_size(0) * (inputs_total + 1);
int shift = n * (inputs_total + 1);
TYPE s = weights[shift + inputs_total];
for(int i = 0; i < inputs_total; i += 4)
s += dot(ToVect4(inputs, i, 1, inputs_total, 0), ToVect4(weights, i, 1, weights_total, shift));
outputs[n] = s;
}
//+------------------------------------------------------------------+
//| Kernel of calculation output gradients |
//| Parameter 'target' vector of target data |
//| 'outputs' vector of previous FF outputs data |
//| 'gradients' vector of gradients |
//| 'loss_function' type of loss function |
//+------------------------------------------------------------------+
__kernel void CalcOutputGradient(__global TYPE *target,
__global TYPE *outputs,
__global TYPE *gradients,
int loss_function)
{
const int n = get_global_id(0);
switch(loss_function)
{
case 0:
gradients[n] = target[n] - outputs[n];
break;
case 1:
gradients[n] = 2 * (target[n] - outputs[n]);
break;
case 2:
gradients[n] = -target[n] / (outputs[n] + 1e-37f) * log(outputs[n] + 1e-37f);
break;
case 3:
gradients[n] = (target[n] - outputs[n]) / (outputs[n] * (outputs[n] - 1) + 1e-37f);
break;
default:
gradients[n] = target[n] - outputs[n];
break;
}
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
__kernel void CalcHiddenGradient(__global TYPE *gradient_inputs,
__global TYPE *weights,
__global TYPE *gradients,
int outputs_total)
{
const int n = get_global_id(0);
const int inputs_total = get_global_size(0);
int weights_total = (inputs_total + 1) * outputs_total;
//---
TYPE grad = 0;
for(int o = 0; o < outputs_total; o += 4)
grad += dot(ToVect4(gradients, o, 1, outputs_total, 0), ToVect4(weights, o, (inputs_total + 1), weights_total, n));
gradient_inputs[n] = grad;
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
__kernel void CalcDeltaWeights(__global TYPE *inputs,
__global TYPE *delta_weights,
__global TYPE *gradients)
{
const int n = get_global_id(0);
const int outputs_total = get_global_size(0);
const int i = get_global_id(1);
const int inputs_total = get_global_size(1);
//---
TYPE grad = gradients[n];
int shift = n * (inputs_total + 1);
delta_weights[shift + i] = inputs[i] * grad + delta_weights[shift + i];
if(i == 0)
delta_weights[shift + inputs_total] += grad;
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
__kernel void SGDUpdate(__global TYPE *delta_weights,
__global TYPE *weights,
int total,
int batch_size,
TYPE learningRate,
TYPE Lambda1,
TYPE Lambda2
)
{
int start = 4 * get_global_id(0);
TYPE4 delta4 = ToVect4(delta_weights, start, 1, total, 0);
TYPE4 weights4 = ToVect4(weights, start, 1, total, 0);
TYPE lr = learningRate / ((TYPE)batch_size);
weights4 -= (TYPE4)(Lambda1) + Lambda2 * weights4;
weights4 += (TYPE4)(lr) * delta4;
D4ToArray(weights, weights4, start, 1, total, 0);
D4ToArray(delta_weights, (TYPE4)(0), start, 1, total, 0);
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
__kernel void MomentumUpdate(__global TYPE *delta_weights,
__global TYPE *weights,
__global TYPE *momentum,
int total, int batch_size,
TYPE learningRate,
TYPE beta,
TYPE Lambda1, TYPE Lambda2)
{
int start = 4 * get_global_id(0);
//---
TYPE4 delta4 = ToVect4(delta_weights, start, 1, total, 0) / ((TYPE4)(batch_size));
TYPE4 weights4 = ToVect4(weights, start, 1, total, 0);
TYPE4 momentum4 = ToVect4(momentum, start, 1, total, 0);
weights4 -= (TYPE4)(Lambda1) + Lambda2 * weights4;
momentum4 = (TYPE4)(learningRate) * delta4 + (TYPE4)(beta) * momentum4;
weights4 += momentum4;
D4ToArray(weights, weights4, start, 1, total, 0);
D4ToArray(momentum, momentum4, start, 1, total, 0);
D4ToArray(delta_weights, (TYPE4)(0), start, 1, total, 0);
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
__kernel void AdaGradUpdate(__global TYPE *delta_weights,
__global TYPE *weights,
__global TYPE *momentum,
int total, int batch_size,
TYPE learningRate,
TYPE Lambda1, TYPE Lambda2)
{
int start = 4 * get_global_id(0);
//---
TYPE4 delta4 = ToVect4(delta_weights, start, 1, total, 0) / ((TYPE4)(batch_size));
TYPE4 weights4 = ToVect4(weights, start, 1, total, 0);
TYPE4 momentum4 = ToVect4(momentum, start, 1, total, 0);
//---
weights4 -= (TYPE4)(Lambda1) + Lambda2 * weights4;
momentum4 = momentum4 + pow(delta4, 2);
weights4 += learningRate / sqrt(momentum4 + 1.0e-37f);
D4ToArray(weights, weights4, start, 1, total, 0);
D4ToArray(momentum, momentum4, start, 1, total, 0);
D4ToArray(delta_weights, (TYPE4)(0), start, 1, total, 0);
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
__kernel void RMSPropUpdate(__global TYPE *delta_weights,
__global TYPE *weights,
__global TYPE *momentum,
int total, int batch_size,
TYPE learningRate,
TYPE beta,
TYPE Lambda1, TYPE Lambda2)
{
int start = 4 * get_global_id(0);
//---
TYPE4 delta4 = ToVect4(delta_weights, start, 1, total, 0) / ((TYPE4)(batch_size));
TYPE4 weights4 = ToVect4(weights, start, 1, total, 0);
TYPE4 momentum4 = ToVect4(momentum, start, 1, total, 0);
//---
weights4 -= (TYPE4)(Lambda1) + Lambda2 * weights4;
momentum4 = beta * momentum4 + (1 - beta) * pow(delta4, 2);
weights4 += delta4 * learningRate / (sqrt(momentum4) + 1.0e-37f);
D4ToArray(weights, weights4, start, 1, total, 0);
D4ToArray(momentum, momentum4, start, 1, total, 0);
D4ToArray(delta_weights, (TYPE4)(0), start, 1, total, 0);
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
__kernel void AdaDeltaUpdate(__global TYPE *delta_weights,
__global TYPE *weights,
__global TYPE *momentumW,
__global TYPE *momentumG,
int total, int batch_size,
TYPE beta1, TYPE beta2,
TYPE Lambda1, TYPE Lambda2)
{
int start = 4 * get_global_id(0);
//---
TYPE4 delta4 = ToVect4(delta_weights, start, 1, total, 0) / ((TYPE4)(batch_size));
TYPE4 weights4 = ToVect4(weights, start, 1, total, 0);
TYPE4 momentumW4 = ToVect4(momentumW, start, 1, total, 0);
TYPE4 momentumG4 = ToVect4(momentumG, start, 1, total, 0);
//---
weights4 -= (TYPE4)(Lambda1) + Lambda2 * weights4;
momentumW4 = beta1 * momentumW4 + (1 - beta1) * pow(weights4, 2);
momentumG4 = beta2 * momentumG4 + (1 - beta2) * pow(delta4, 2);
weights4 += delta4 * sqrt(momentumW4) / (sqrt(momentumG4) + 1.0e-37f);
D4ToArray(weights, weights4, start, 1, total, 0);
D4ToArray(momentumW, momentumW4, start, 1, total, 0);
D4ToArray(momentumG, momentumG4, start, 1, total, 0);
D4ToArray(delta_weights, (TYPE4)(0), start, 1, total, 0);
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
__kernel void AdamUpdate(__global TYPE *delta_weights,
__global TYPE *weights,
__global TYPE *momentumM,
__global TYPE *momentumV,
int total, int batch_size,
TYPE learningRate,
TYPE beta1, TYPE beta2,
TYPE Lambda1, TYPE Lambda2)
{
int start = 4 * get_global_id(0);
//---
TYPE4 delta4 = ToVect4(delta_weights, start, 1, total, 0) / ((TYPE4)(batch_size));
TYPE4 weights4 = ToVect4(weights, start, 1, total, 0);
TYPE4 momentumM4 = ToVect4(momentumM, start, 1, total, 0);
TYPE4 momentumV4 = ToVect4(momentumV, start, 1, total, 0);
//---
momentumM4 = beta1 * momentumM4 + (1 - beta1) * delta4;
momentumV4 = beta2 * momentumV4 + (1 - beta2) * pow(delta4, 2);
TYPE4 m = momentumM4 / (1 - beta1);
TYPE4 v = momentumV4 / (1 - beta2);
weights4 -= (TYPE4)(Lambda1) + Lambda2 * weights4;
weights4 += learningRate * m / (sqrt(v) + 1.0e-37f);
D4ToArray(weights, weights4, start, 1, total, 0);
D4ToArray(momentumM, momentumM4, start, 1, total, 0);
D4ToArray(momentumV, momentumV4, start, 1, total, 0);
D4ToArray(delta_weights, (TYPE4)(0), start, 1, total, 0);
}
//+------------------------------------------------------------------+
//| Feed-forward kernel of the convolutional layer |
//| Parameters: 'inputs' input data vector |
//| 'weights' weight matrix |
//| 'outputs' vector of results |
//| 'inputs_total' size of input data vector |
//| 'window' size of input data analysis window |
//| 'step' window moving step |
//| 'window_out' number of filters |
//+------------------------------------------------------------------+
__kernel void ConvolutionFeedForward(__global TYPE *inputs,
__global TYPE *weights,
__global TYPE *outputs,
int inputs_total,
int window,
int step,
int window_out,
int transposed_out)
{
const int n = get_global_id(0);
const int neurons = get_global_size(0);
const int weights_total = (window + 1) * window_out;
int shift = n * step;
for(int w = 0; w < window_out; w++)
{
int out = (transposed_out == 1 ? w + n * window_out : w * neurons + n);
int shift_weights = w * (window + 1) ;
if((shift_weights + window) >= weights_total)
break;
TYPE s = weights[shift_weights + window];
for(int i = 0; i < window; i += 4)
s += dot(ToVect4(inputs, i, 1, inputs_total, shift),
ToVect4(weights, i, 1, shift_weights + window, shift_weights));
outputs[out] = s;
}
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
__kernel void ConvolutionCalcHiddenGradient(__global TYPE *gradient_inputs,
__global TYPE *weights,
__global TYPE *gradients,
int window,
int step,
int window_out,
int neurons,
int transposed_out)
{
const int n = get_global_id(0);
const int inputs_total = get_global_size(0);
int weights_total = (window + 1) * window_out;
//---
TYPE grad = 0;
int w_start = n % step;
int r_start = max((n - window + step) / step, 0);
int total = (window - w_start + step - 1) / step;
total = min((n + step) / step, total);
for(int i = 0; i < total; i ++)
{
int row = r_start + i;
if(row >= neurons)
break;
for(int wo = 0; wo < window_out; wo++)
{
int shift_g = (transposed_out == 1 ? row * window_out + wo : row + wo * neurons);
int shift_w = w_start + (total - i - 1) * step + wo * (window + 1);
grad += gradients[shift_g] * weights[shift_w];
}
}
gradient_inputs[n] = grad;
}
//+------------------------------------------------------------------+
//| |
//+------------------------------------------------------------------+
__kernel void ConvolutionCalcDeltaWeights(__global TYPE *inputs,
__global TYPE *delta_weights,
__global TYPE *gradients,
int inputs_total,
int step,
int neurons,
int transposed_out)
{
const int inp_w = get_global_id(0);
const int w = get_global_id(1);
const int window = get_global_size(0) - 1;
const int window_out = get_global_size(1);
//---
int shift_delt = w * (window + 1) + inp_w;
TYPE value = 0;
if(inp_w == window)
{
for(int n = 0; n < neurons; n ++)
value += gradients[transposed_out == 1 ? w + n*window_out : w * neurons + n];
}
else
for(int n = 0; n < neurons; n ++)
{
int shift_inp = n * step + inp_w;
if(shift_inp >= inputs_total)
break;
value += inputs[shift_inp] * gradients[transposed_out == 1 ? w + n*window_out : w * neurons + n];
}
delta_weights[shift_delt] += value;
}
//+------------------------------------------------------------------+
//| Feed-forward kernel of the pooling layer |
//| Parameters: 'inputs' input data vector |
//| 'outputs' vector of results |
//| 'inputs_total' size of input data vector |
//| 'input_neurons' vector size of 1st input data filter |
//| 'window' size of input data analysis window |
//| 'step' window moving step |
//| 'activation' type of activation function |
//+------------------------------------------------------------------+
__kernel void ProofFeedForward(__global TYPE *inputs,
__global TYPE *outputs,
int inputs_total,
int input_neurons,
int window,
int step,
int activation)
{
const int n = get_global_id(0);
const int w = get_global_id(1);
const int neurons = get_global_size(0);
const int window_out = get_global_size(1);
int shift = n * step;
int out = w * neurons + n;
int shift_inp = w * input_neurons;
TYPE s = 0;
TYPE k = (TYPE)1 / (TYPE)window;
TYPE4 k4 = (TYPE4)(k);
for(int i = 0; i < window; i += 4)
switch(activation)
{
case 0:
s += dot(ToVect4(inputs, i, 1, min(shift_inp + input_neurons, inputs_total), shift_inp + shift),
k4);
break;
case 1:
s = Max4(ToVect4(inputs, i, 1, min(shift_inp + input_neurons, inputs_total), shift_inp + shift), s);
break;
default:
break;
}
outputs[out] = s;
}
//+------------------------------------------------------------------+
//| Backpropagation kernel of the pooling layer |
//| Parameters: 'inputs' input data vector |
//| 'gradient_inputs' previous layer gradients vector |
//| 'outputs' vector of results |
//| 'gradients' current layer gradients vector |
//| 'inputs_total' size of input data vector |
//| 'outputs_total' size of outputs vector |
//| 'window' size of input data analysis window |
//| 'step' window moving step |
//| 'neurons' vector size if 1st filter of outputs |
//| 'activation' type of activation function |
//+------------------------------------------------------------------+
__kernel void ProofCalcHiddenGradient(__global TYPE *inputs,
__global TYPE *gradient_inputs,
__global TYPE *outputs,
__global TYPE *gradients,
int inputs_total,
int outputs_total,
int window,
int step,
int neurons,
int activation)
{
const int n = get_global_id(0);
const int w = get_global_id(1);
const int input_neurons = get_global_size(0);
const int window_out = get_global_size(1);
//---
int start = max((n - window + step) / step, 0);
int stop = min((n + step - 1) / step + 1, neurons);
TYPE grad = 0;
int shift_inp = w * input_neurons + n;
if(shift_inp >= inputs_total)
return;
TYPE inp = inputs[shift_inp];
int shift_out = w * neurons;
for(int o = start; o < stop; o ++)
{
int shift_g = shift_out + o;
if(shift_g >= outputs_total)
break;
switch(activation)
{
case 0:
grad += gradients[shift_g] / (TYPE)window;
break;
case 1:
grad += (outputs[shift_g] == inp ? gradients[shift_g] : 0);
break;
default:
break;
}
}
gradient_inputs[shift_inp] = grad;
}
//+------------------------------------------------------------------+
//| LSTM block feed-forward kernel |
//| Parameters: 'forgetgate' forget gate |
//| 'inputgate' input gate |
//| 'outputgate' output gate |
//| 'newcontent' new content |
//| 'memory' memory stream |
Copy permalink