Что под капотом у нейронной сети. Нейросеть c точки зрения математики и программирования / forpes.ru

Главная
Что под капотом у нейронной сети. Нейросеть c точки зрения математики и программирования

Что под капотом у нейронной сети. Нейросеть c точки зрения математики и программирования +8

18.06.2023 05:35

aleksandr__sokol 26 15000 Источник

Здравствуйте, меня зовут Александр, я backend-разработчик.

Часто искусственные нейронные сети рассматриваются или с точки зрения математических моделей, или с точки зрения написания программ на конкретном языке. Как в притче о слоне и слепых мудрецах.

Цель данной публикации – комплексное рассмотрение строения искусственных нейронных сетей c точки зрения и математики и программного кода. В данной работе нейронная сеть реализуется на языке Python с использованием библиотеки tensorflow.keras. Статья сосредоточена в основном на строении и функционировании искусственной нейронной сети, поэтому такие этапы как обучение и т.д. в ней не затрагиваются.

Несмотря на большое разнообразие вариантов нейронных сетей, все они имеют общие черты. Так, все они, так же, как и мозг человека, состоят из большого числа связанных между собой однотипных элементов – нейронов, которые имитируют нейроны головного мозга. С точки зрения математики, в нейронной сети каждый нейрон представляет из себя функцию, зависящую от значений входных сигналов, весов этих сигналов и некой активационной функции.

Прежде всего необходимо остановиться на понятии активационной функции. Она может быть представлена в виде любой функцией от нескольких аргументов, возвращающей одно значение. Данная функция нужна, чтобы определить при каких входных значениях должен включаться (активироваться) нейрон, т.е. проводить сигнал. А при каких нет. Чаще всего используются следующие функции:

Линейная функция активации который имеет вид
Сигмоид который имеет вид ${f(x,a)} = {1\over 1 + e^{-ax}}$
ReLu вида $f(x,a) = \begin{cases} x & \quad \text{x > 0}\\ ax & \quad \text{x <= 0} \end{cases}$

В расчетах будем использовать активационную функцию ReLu с параметром a = 1. Т.к. активационная функция имеет вид f(x) = x, для положительных значений параметров ей можно пренебречь.

Вернемся к строению нейрона. На каждый вход нейрона подаются значения (X), которые затем распространяются по межнейронным связям (синапсисам). У синапсов есть один параметр — W (вес), благодаря которому входная информация изменяется при переходе от одного нейрона к другому.

X1 … Xn – значение входных сигналов, Y – выходной сигнал, W1…Wn – вес определяет, насколько соответствующий вход нейрона влияет на его состояние.

Таким образом нейрон так же представляет из себя некую функцию вида

$S_n(X,W) = f({\sum_{i=1}^n(X_iW_i)})$

где f – активационная функция.

Таким образом, искусственная нейронная сеть – это линейный многочлен со множеством параметров.

Рассмотрим пример. Исследуется зависимость Y от двух параметров X1 и X2.

В ходе неких экспериментов получены следующие зависимости.

X1	X2	Y
1	1	2
2	1	3
3	1	4
4	1	5
5	1	6
1	2	3
2	2	4
3	2	5
4	2	6
5	2	7
1	5	6

Не сложно заметить, что зависимость
Но допустим, мы этого не знаем. Попробуем найти значения Y обучив нейронную сеть.

Мы построим несколько видов нейросетей (с разным числом нейронов), а затем с помощью библиотеки keras в python найдем нужные веса. Для расчетов будет использоваться многослойная нейронная сеть.

from tensorflow.keras.models import Sequential
model = Sequential()

Для нахождения весов на каждом слое используется команда

for layer in model.layers: 
  print(layer.get_weights())

Рассмотрим две схемы нейронной сети:

В первом случае нейронная сеть имеет два слоя, на каждом слое по одному нейрону (вычисленный вес показан на рисунке).

model.add(Dense(1, input_shape=X.shape[1:], kernel_constraint=non_neg()))
model.add(layers.Activation(keras.activations.relu))
model.add(Dense(1, activation='linear', kernel_constraint=non_neg()))

Выполним проверку, т.е. вычислим многочлен Y = (X1 * W11 + X2 * W12) * W21

Для получения аналогичных данных используется метод predict

model.predict(…)

X1	X2	Y = X1 + X2	Вычисленный результат
1	1	2	(10.7717732 + 10.7717735)*1.295717 = 2
10	10	20	(100.7717732 + 100.7717735)*1.295717 = 20
5.5	2.5	8	(5.50.7717732 + 2.50.7717735)*1.295717 = 8
10.3	2.1	12.4	(10.30.7717732 + 2.10.7717735)*1.295717 = 12.4
8.3	20.1	28.4	(8.30.7717732 + 20.10.7717735)*1.295717 = 28.4

Во втором случае нейронная сеть имеет два слоя, на первом слое два нейрона, а на втором они нейрон (вычисленный вес показан на рисунке).

model.add(Dense(1, input_shape=X.shape[1:], kernel_constraint=non_neg()))
model.add(layers.Activation(keras.activations.relu))
model.add(Dense(1, activation='linear', kernel_constraint=non_neg()))

Вычислим многочлен Y = (X1 * W11 + X2 * W12) * W21 + (X1 * W13 + X2 * W14) * W22

X1	X2	Y = X1 + X2	Вычисленный результат
1	1	2	(10.75749624 + 10.7609735)* 0.67330104 + (10.6390331+ 1 0.88485044)*0.6438818 = 2
10	10	20	(10* 0.75749624 + 10* 0.7609735)0.67330104 + (10 0.6390331+ 10* 0.88485044)*0.6438818 = 20
5.5	2.5	8	(5.5* 0.75749624 + 2.5* 0.7609735)0.67330104 + (5.5 0.6390331+ 2.5* 0.88485044)*0.6438818 = 8
10.3	2.1	12.4	(10.3* 0.75749624 + 2.1* 0.7609735)0.67330104 + (10.3 0.6390331+ 2.1* 0.88485044)*0.6438818 = 11.76
8.3	20.1	28.4	(8.3* 0.75749624 + 20.1* 0.7609735)0.67330104 + (8.3 0.6390331+ 20.1* 0.88485044)*0.6438818 = 29.4

Итак, мы рассмотрели строение нейрона и простейших нейронных сетей, которые, с точки зрения математики, представляют из себя многочлены с большим числом параметров (весов). Обучение нейронной сети как раз представляет из себя нахождение этих параметров.

Комментарии (26)

da-nie
18.06.2023 05:58
#25662258
+3
Это очень здорово, но гораздо интереснее было бы рассказать, как именно происходит обучение пакетами на GPU.
1) Весь пакет прогоняется на GPU параллельно? Обычно, столько памяти нет.
2) На CPU сделан цикл по элементам пакета. Один элемент копируется на GPU и производится обучение на нём с сумированием весов. В этом случае непонятно, как делать нормализацию?
3) На GPU копируется пакет и один слой сети. Выполняется проход по одному слою сети для всех элементов пакета. Далее, пакет выбрасывается и выход слоя является новым пакетом для следующего слоя. Памяти нужно меньше, чем в пункте 1, но все промежуточные результаты придётся копировать на CPU (они нужны для обратного прохода). А такое копирование очень медленное.
1. da-nie
  18.06.2023 05:58
  #25662842
  +7
  Судя по минусаторам, все они, конечно, знают, как же сделано на практике в пакетах типа keras и tensorflow? Не так ли? ;) Но при этом они не знают по какой математике считается нейрон, раз аплодируют этой абсолютно примитивной статье?
1. iskateli
  18.06.2023 05:58
  #25663652
  интереснее было бы рассказать, как именно происходит обучение
  
  поэтому такие этапы как обучение и т.д. в ней не затрагиваются
  интересно это было написано до вашего комментария? но в любом случае напрасно опускать такой важный этап как обучение, раз "Цель данной публикации – комплексное рассмотрение "
  1. da-nie
    18.06.2023 05:58
    #25663756
    Нет, это было изначально написано.
    Потому я и написал, что интересно было бы прочесть вот про указанное (а не одно и то же в каждой статье по нейронкам).

gleb_l
18.06.2023 05:58
#25662330
+4
Нейрон - это взвешивающий сумматор с (опционально) нелинейным выходом. Хорошо, если функция активации дифференцируема на R. Тогда его можно учить методом обратного распространения ошибки. Если нет - таким способом учить его невозможно (только Больцманом - закреплением случайных флуктуаций весов по локально-лучшему результату)

fedorro

18.06.2023 05:58

#25662840

+11

Тут автору каждой сотой статьи про то как работают нейронные сети ачивку выдают, или в чем смысл их постоянно постить?)

da-nie

18.06.2023 05:58

#25662850

+11

Каждый автор считает своим долгом донести внезапно полученные знания о нейросетях (часто очень примитивные). :) Жаль только, что в математике и принципах обучения они редко что-то понимают, ибо популярные пакеты всё это тщательно скрывают. В результате по тем же свёрточным сетям хрен найдёшь как их обучать. Или по сетям GAN как считать ошибку выходного слоя дискриминатора. Или как делать нормализацию на пальцах. Или как же обучать пакетом на GPU. Или как сделать быструю свёртку, обратную свёртку и коррекцию коэффициентов на GPU. А в остальном да, все специалисты. :)

А здесь статья с очень громким названием "Что под капотом у нейронной сети. Нейросеть c точки зрения математики и программирования". То есть, ожидается как минимум вся подноготная. Но… имеем то, что имеем. Мой вон вопрос-предложение выше вызывает у многих баттхерт. А я всего лишь спросил ту информацию, которую я не сумел пока ещё нигде найти.

fedorro

18.06.2023 05:58

#25663028

Информацию в смысле как на GPU много данных обсчитывается? Если да - то ответ прост - батчами, т.е. не все примеры сразу загоняются в память, а столько, сколько влазит, после каждого прогона батча обновляются веса. Да, при этом градиент считается не по всему обучающему множеству сразу, но зато в память влазит. Мне кажется во всех современных нейросетевых библиотеках есть параметр batch_size, который надо подгонять под то чтобы побольше памяти за раз использовать, но не поймать Out Of Memory. Как то так.

da-nie

18.06.2023 05:58

#25663070

Нет, ответ не прост. Даже для достаточно примитивных сетей требуется весьма много памяти, и пакет у вас будет очень маленьким (и вообще выродится в один образ). Потом, от размера пакета будет зависеть нормализация.
Вот тут я делал на С++ такую сеть с самодельной работой с GPU.

Так вот, в мою карту 1060 с 3 ГБ поместилось за раз не более 100 входных изображений из 60000. Для более сложной сети вообще помещается один образ и не более. И как нормализацию делать в таком случае? Вот именно.

fedorro

18.06.2023 05:58

#25663086

Ну 100 - это уже хорошо, на то они и мини-батчи. А один пример за раз - тоже не новость, называется "online learning". Вот тут сравнение всех трех вариантов с примерами и т.п.: Full batch, mini-batch, and online learning

da-nie
18.06.2023 05:58
#25663114
+1
Так 100 — это для очень-очень простой сети. Когда я 160x120 генерировать пытался на свёрточной сети, там вообще штук 5 образов влезало за раз. И всю память жрут полносвязаные слои (матрицы весов быстро нарастают). Сами ядра свёрток мало жрут.
Поэтому мне вот кажется, вероятно, послойно там всё обрабатывается внутри с выбросом использованных слоёв.

Есть ещё вопросы к быстрой свёртке. Вот есть алгоритм свёртки через умножение. Я его сделал влоб сперва. Какой наивный! :) Одна (!) матрица такой свёртки для картинки 300x300 весила больше 8 ГБ. Ладно, переделал с генерацией матрицы на лету и умножением матрицы с shared-памятью. Заработало. Но скорость упала в 100 раз по сравнению с простой распараллеленой свёрткой без shared-памяти на GPU — оно и понятно, очень много лишних умножений. Мда.

da-nie

18.06.2023 05:58

#25663136

Кстати! Вот смотрите, есть у вас CUDA. Там у вас есть трёхмерные блоки и трёхмерные потоки внутри блоков. Смотрите какая штука вас ожидает.
Быстрое матричное умножение (а оно потребуется для полносвязных сетей) требует задать в потоках блок x=16, y=16,z=1 (число всех потоков внутри блока сильно ограничено! 512, 1024 и, может, уже чуть больше на новых картах. То есть, размерность полной матрицы вы в него не вставите никогда для матриц превышающих по размеру число потоков). Осталась размерность по Z, но что там поместится? 1024/(256)=4. Вот максимум размерности по Z. Вообще не используем её в этом случае. Как видите, мы уже заняли размерность потока. Теперь берём блоки. В x и y мы помещаем номер блока 16x16 (который обрабатывают потоки) по x и по y внутри матрицы результата умножения. А что в z? Номер батча? Ну уж нет. У нас не матрицы, а 3D-тензоры в общем случае, поэтому там z для тензора.
То есть, мы уже заняли всю параллельную работу на видяхе на одну только матрицу. :)

fedorro

18.06.2023 05:58

#25663190

Ну если так микроскопом гвозди забивать - конечно никаких ресурсов не хватит ????Откуда матрицы свертки на 8 Gb? Обычно они по сотне байтов занимают (зависит от размера, конечно). Я на MNIST первую сверточную сеть запускал LeNet или AlexNet - не помню, но она обучилась нормально на двухядерном процессоре с 4Gb памяти, без всяких GPU. Посмотрели бы как это сделано в том же TensorFlow - он на C++ написан, а на питоне только биндинги для удобного взаимодействия (а Keras - это ещё немного абстракций на питоне для того же TensorFlow, не более). PyTorch - так же, как и numpy и т.д.

da-nie

18.06.2023 05:58

#25663210

-1

Вот откуда:

Или так нагляднее:

Имея быстрое умножение матриц на GPU (на shared-памяти) очень хочется его использовать для ускорения свётки (так как shared-память я пока использовать для произвольных параметров свёртки не придумал как).
И вот есть статья.

Посмотрели бы как это сделано в том же TensorFlow

Это не так-то просто будет понять.

vassabi
18.06.2023 05:58
#25663436
+1
ого, вот это вы мощно - это ж сколько движений для создания матриц и перемножения!

а вы не пробовали

1) сделать из вектора f вектор ff добавив в него два нуля в начале и два нуля в конце

2) умножить вектор ff три раза на компоненты вектора g (получив три вектора - ffg0, ffg1 и ffg2)
3) проссуммировать компоненты этих векторов по диагонали (т.е. взяв ffg0[i+2] + ffg0[i+1] + ffg0[i])
и все - никаких ветвлений и лишних движений.

PS: еще можно поиграться с тем, чтобы 2 и 3 делать в одном и том же массиве (держать значения ffg0, ffg1 и ffg2 с шагом 3, т.е. ffg0[0] -> ffg[0], ffg1[0] -> ffg[1],ffg2[0] -> ffg[2], ffg0[1] -> ffg[3] , ffg1[1] -> ffg[4] ...)
1. da-nie
  18.06.2023 05:58
  #25663634
  это ж сколько движений для создания матриц и перемножения!
  
  Это только вариант на аналогах всяких матриц Теплица. Оказался неудачным.
  
  а вы не пробовали
  
  Я пока не очень понял, что вы имеете в виду. В исходном варианте и g и f не вектора, а матрицы. Их напрямую так не перемножишь.
  Я сейчас пробую понять, что сделано тут.
  1. vassabi
    18.06.2023 05:58
    #25663728
    насколько я понял, там делают предрасчет сумм окна свертки (ибо это можно сделать одним проходом, очень эффективно), а потом уже умножают на коэфициенты
  1. da-nie
    18.06.2023 05:58
    #25663760
    Меня смущает вот что:
    
    Исходя из того, что K = srcC kernelY kernelX / group, эффективность метода особенно низка для входных свёрточных слоев. А для depthwise convolution матричный метод вообще проигрывает тривиальной реализации.
    
    Дело ещё в том, что кроме прямой и обратной свётки (и перевёрнутой) есть ещё расчёт поправок коэффициентов ядер. Этот этап тоже надо думать как ускорить.

fedorro

18.06.2023 05:58

#25663834

Ну и какая-же это оптимизация, если вместо маленькой матрицы, которую можно вообще в константы шейдера задать, у вас 8Gb чисел ...

Задать свертку шейдером и запустить на просчет на каждом участке видеокарты - оно как раз распараллелится на кол-во CUDA-ядер, и вся свертка в два присеста посчитается, а памяти надо будет порядка размера обрабатываемого изображения.

da-nie

18.06.2023 05:58

#25663868

Как я уже раза два написал, это вариант через умножение матриц.

Задать свертку шейдером

Не использовал шейдеры, но, полагаю, максимального быстродействия вы не получите, если просто сделаете свёртку шейдером.

Вот смотрите, как сделаны у меня свёртки и как они распараллелены на блоки и потоки для 3D-тензоров.

#ifndef C_TENSOR_CONV_H
#define C_TENSOR_CONV_H

//****************************************************************************************************
//Операции свёртки над тензорами произвольной размерности
//****************************************************************************************************

//****************************************************************************************************
//подключаемые библиотеки
//****************************************************************************************************
#include "ctensor.cu.h"
#include "ctensormath.cu.h"
#include <vector>

//****************************************************************************************************
//макроопределения
//****************************************************************************************************

//****************************************************************************************************
//константы
//****************************************************************************************************

//****************************************************************************************************
//предварительные объявления
//****************************************************************************************************

//****************************************************************************************************
//прототипы функций
//****************************************************************************************************

//****************************************************************************************************
///!Операции свёртки над тензорами произвольной размерности
//****************************************************************************************************

template<class type_t>
class CTensorConv
{
 public:
  //-перечисления---------------------------------------------------------------------------------------
  //-структуры------------------------------------------------------------------------------------------
  //-константы------------------------------------------------------------------------------------------
 private:
  //-переменные-----------------------------------------------------------------------------------------
 public:
  //-конструктор----------------------------------------------------------------------------------------
  //-конструктор копирования----------------------------------------------------------------------------
  //-деструктор-----------------------------------------------------------------------------------------
 public:
  //-открытые функции-----------------------------------------------------------------------------------
  static void ForwardConvolution(CTensor<type_t> &cTensor_Output,const CTensor<type_t> &cTensor_Image,const std::vector<CTensor<type_t> > &cTensor_Kernel,const std::vector<type_t> &bias,uint32_t step_y,uint32_t step_x,uint32_t padding_y,uint32_t padding_x);///<прямая свёртка
  static void BackwardConvolution(CTensor<type_t> &cTensor_OutputDelta,const CTensor<type_t> &cTensor_Delta,const std::vector<CTensor<type_t> > &cTensor_Kernel,const std::vector<type_t> &bias);///<обратная свёртка
  static void CreateDeltaWeightAndBias(std::vector<CTensor<type_t> > &cTensor_dKernel,std::vector<type_t> &dbias,const CTensor<type_t> &cTensor_Image,const CTensor<type_t> &cTensor_Delta);///<вычисление поправок весов и смещений
  static void CreateBackDeltaWeightAndBias(std::vector<CTensor<type_t> > &cTensor_dKernel,std::vector<type_t> &dbias,const CTensor<type_t> &cTensor_Image,const CTensor<type_t> &cTensor_Delta);///<создание поправок весов и смещений для обратной свёртки
 private:
  //-закрытые функции-----------------------------------------------------------------------------------
};

//****************************************************************************************************
//конструктор и деструктор
//****************************************************************************************************

//****************************************************************************************************
//закрытые функции
//****************************************************************************************************

//----------------------------------------------------------------------------------------------------
//
//----------------------------------------------------------------------------------------------------

//****************************************************************************************************
//открытые функции
//****************************************************************************************************

//****************************************************************************************************
//статические функции
//****************************************************************************************************

//----------------------------------------------------------------------------------------------------
//функция CUDA для выполнения прямой свёртки
//----------------------------------------------------------------------------------------------------
template<class type_t>
__global__ void CUDAForwardConvolutionFunction(STensorKernel<type_t> tensor_output,STensorKernel<type_t> tensor_image,STensorKernel<type_t> tensor_kernel,type_t** kernel_item_ptr_array,type_t* bias_ptr,int32_t padding_x,int32_t padding_y,int32_t step_x,int32_t step_y)
{
 size_t blockCol=blockIdx.x;
 size_t blockRow=blockIdx.y;
 size_t z=blockIdx.z;
 //координаты элементов блока в выходном тензоре
 size_t x=threadIdx.x;
 size_t y=threadIdx.y;

 x+=blockCol*CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE;
 y+=blockRow*CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE;

 int32_t output_z=z;
 int32_t output_x=tensor_output.GetSizeX();
 int32_t output_y=tensor_output.GetSizeY();

 if (x>=output_x) return;
 if (y>=output_y) return;

 int32_t kernel_x=tensor_kernel.GetSizeX();
 int32_t kernel_y=tensor_kernel.GetSizeY();
 int32_t kernel_z=tensor_kernel.GetSizeZ();

 int32_t input_x=tensor_image.GetSizeX();
 int32_t input_y=tensor_image.GetSizeY();

 //настроим ядро
 tensor_kernel.TensorData_Ptr=kernel_item_ptr_array[output_z];

 type_t sum=bias_ptr[output_z];//сразу прибавляем смещение
 //применяем фильтр
 for(int32_t ky=0;ky<kernel_y;ky++)
 {
  int32_t y0=step_y*y+ky-padding_y;
  if (y0<0 || y0>=input_y) continue;
  for(int32_t kx=0;kx<kernel_x;kx++)
  {
   int32_t x0=step_x*x+kx-padding_x;
   //игнорируем элементы вне границ входного тензора
   if (x0<0 || x0>=input_x) continue;
   //проходимся по всей глубине тензора и считаем сумму
   for(int32_t z=0;z<kernel_z;z++)
   {
    type_t kernel=tensor_kernel.GetElement(z,ky,kx);
    type_t image=tensor_image.GetElement(z,y0,x0);
    sum+=kernel*image;
   }
  }
 }
 tensor_output.SetElement(output_z,y,x,sum);

/*
 int32_t kx=threadIdx.x;
 int32_t ky=threadIdx.y;
 int32_t kz=threadIdx.z;

 int32_t x=blockIdx.x;
 int32_t y=blockIdx.y;

 //применяем фильтр
 int32_t y0=step_y*y+ky-padding_y;
 int32_t x0=step_x*x+kx-padding_x;
 //игнорируем элементы вне границ входного тензора
 if (y0<0 || y0>=tensor_image.GetSizeY() || x0<0 || x0>=tensor_image.GetSizeX())
 {
  summ[kz][ky][kx]=0;
 }
 else
 {
  type_t kernel=tensor_kernel.GetElement(kz,ky,kx);
  type_t image=tensor_image.GetElement(kz,y0,x0);
  summ[kz][ky][kx]=kernel*image;
 }

 __syncthreads();

 //делаем суммирование результата свёртки
 if (kx==0 && ky==0 && kz==0)
 {
  type_t s=0;
  for(int32_t klz=0;klz<tensor_kernel.GetSizeZ();klz++)
  {
   for(int32_t kly=0;kly<tensor_kernel.GetSizeY();kly++)
   {
    for(int32_t klx=0;klx<tensor_kernel.GetSizeX();klx++)
    {
     s+=summ[klz][kly][klx];
    }
   }
  }
  tensor_output.SetElement(output_z,y,x,s);
 }
 __syncthreads();
 */
}

//----------------------------------------------------------------------------------------------------
/*!прямая свёртка
*/
//----------------------------------------------------------------------------------------------------
template<class type_t>
void CTensorConv<type_t>::ForwardConvolution(CTensor<type_t> &cTensor_Output,const CTensor<type_t> &cTensor_Image,const std::vector<CTensor<type_t> > &cTensor_Kernel,const std::vector<type_t> &bias,uint32_t step_y,uint32_t step_x,uint32_t padding_y,uint32_t padding_x)
{
 //вычисляем размеры выходного тензора
 int32_t output_z=cTensor_Kernel.size();
 if (output_z==0) throw("Для прямой свёртки требуется хотя бы одно ядро свёртки");
 if (output_z!=bias.size()) throw("Для прямой свёртки требуется чтобы количество ядер и смещений совпадало");

 int32_t kernel_x=cTensor_Kernel[0].Size_X;
 int32_t kernel_y=cTensor_Kernel[0].Size_Y;
 int32_t kernel_z=cTensor_Kernel[0].Size_Z;

 int32_t input_y=cTensor_Image.Size_Y;
 int32_t input_x=cTensor_Image.Size_X;
 int32_t input_z=cTensor_Image.Size_Z;

 int32_t output_x=(input_x-kernel_x+2*padding_x)/step_x+1;
 int32_t output_y=(input_y-kernel_y+2*padding_y)/step_y+1;

 if (cTensor_Output.Size_X!=output_x || cTensor_Output.Size_Y!=output_y || cTensor_Output.Size_Z!=output_z) throw("Ошибочная размерность выходного тензора для свёртки");

 if (input_z!=kernel_z) throw("Для прямой свёртки требуется чтобы глубина фильтров и входного тензора совпадали");

 STensorKernel<type_t> sTensorKernel_Output(cTensor_Output);//выходной тензор
 STensorKernel<type_t> sTensorKernel_Image(cTensor_Image);//входной тензор

 cTensor_Image.CopyToDevice();

 //копируем на видеокарту указатели и смещения
 CCUDADeviceVector<type_t> bias_array(bias.size());
 bias_array.copy_host_to_device(&bias[0],bias.size());

 CCUDADeviceVector<type_t *> kernel_item_ptr_array(cTensor_Kernel.size());
 std::vector<type_t *> item_ptr(cTensor_Kernel.size());
 for(size_t n=0;n<cTensor_Kernel.size();n++)
 {
  cTensor_Kernel[n].CopyToDevice();
  item_ptr[n]=cTensor_Kernel[n].GetDeviceVector().get();
 }
 kernel_item_ptr_array.copy_host_to_device(&item_ptr[0],item_ptr.size());
 //выполняем свёртку
 STensorKernel<type_t> sTensorKernel_Kernel(cTensor_Kernel[0]);

 dim3 thread(CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE,CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE);

 size_t block_x=output_x/thread.x;
 if (output_x%thread.x) block_x++;
 size_t block_y=output_y/thread.y;
 if (output_y%thread.y) block_y++;
 size_t block_z=cTensor_Kernel.size();

 dim3 blocks(block_x,block_y,block_z);
 if (blocks.x==0) blocks.x=1;
 if (blocks.y==0) blocks.y=1;
 if (blocks.z==0) blocks.z=1;
 CUDAForwardConvolutionFunction<type_t><<<blocks,thread>>>(sTensorKernel_Output,sTensorKernel_Image,sTensorKernel_Kernel,kernel_item_ptr_array.get(),bias_array.get(),padding_x,padding_y,step_x,step_y);
 HANDLE_ERROR(cudaGetLastError());
 HANDLE_ERROR(cudaDeviceSynchronize());

/*
 dim3 thread(cTensor_Kernel.size(),1,1);
 dim3 blocks(output_x,output_y,1);
 if (blocks.x==0) blocks.x=1;
 if (blocks.y==0) blocks.y=1;
 if (blocks.z==0) blocks.z=1;
 CUDAForwardConvolutionFunction<type_t><<<blocks,thread>>>(sTensorKernel_Output,sTensorKernel_Image,sTensorKernel_Kernel,kernel_item_ptr_array.get(),bias_array.get(),padding_x,padding_y,step_x,step_y);
 HANDLE_ERROR(cudaGetLastError());
 HANDLE_ERROR(cudaDeviceSynchronize());
 */

 cTensor_Output.SetDeviceOnChange();

}

//----------------------------------------------------------------------------------------------------
//функция CUDA для выполнения обратной свёртки
//----------------------------------------------------------------------------------------------------
template<class type_t>
__global__ void CUDABackwardConvolutionFunction(STensorKernel<type_t> tensor_output_delta,STensorKernel<type_t> tensor_delta,STensorKernel<type_t> tensor_kernel,type_t** kernel_item_ptr_array,size_t kernel_amount,type_t* bias_ptr,int32_t padding_x,int32_t padding_y,int32_t step_x,int32_t step_y)
{
 size_t blockCol=blockIdx.x;
 size_t blockRow=blockIdx.y;
 size_t z=blockIdx.z;
 //координаты элементов блока в выходном тензоре
 size_t x=threadIdx.x;
 size_t y=threadIdx.y;

 x+=blockCol*CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE;
 y+=blockRow*CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE;

 int32_t output_z=z;
 int32_t output_x=tensor_output_delta.GetSizeX();
 int32_t output_y=tensor_output_delta.GetSizeY();

 if (x>=output_x) return;
 if (y>=output_y) return;
/*

 int32_t x=blockIdx.x;
 int32_t y=blockIdx.y;

 int32_t z=threadIdx.x;

 int32_t output_z=z;
 */

 int32_t kernel_x=tensor_kernel.GetSizeX();
 int32_t kernel_y=tensor_kernel.GetSizeY();
 int32_t kernel_z=tensor_kernel.GetSizeZ();

 int32_t input_x=tensor_delta.GetSizeX();
 int32_t input_y=tensor_delta.GetSizeY();

 type_t sum=0;
 //применяем фильтр
 for(int32_t ky=0;ky<kernel_y;ky++)
 {
  int32_t y0=y+ky-padding_y;
  if (y0<0 || y0>=input_y) continue;
  for(int32_t kx=0;kx<kernel_x;kx++)
  {
   int32_t x0=x+kx-padding_x;
   //игнорируем элементы вне границ входного тензора
   if (x0<0 || x0>=input_x) continue;
   //проходимся по всей глубине тензора и считаем сумму
   for(int32_t k=0;k<kernel_amount;k++)
   {
    //настроим ядро
    tensor_kernel.TensorData_Ptr=kernel_item_ptr_array[k];
    //считаем свёртку
    type_t kernel=tensor_kernel.GetElement(output_z,kernel_y-1-ky,kernel_x-1-kx);
    type_t delta=tensor_delta.GetElement(k,y0,x0);

    sum+=kernel*delta;
   }
  }
 }
 tensor_output_delta.SetElement(output_z,y,x,sum);
}

//----------------------------------------------------------------------------------------------------
/*!обратная свёртка
*/
//----------------------------------------------------------------------------------------------------
template<class type_t>
void CTensorConv<type_t>::BackwardConvolution(CTensor<type_t> &cTensor_OutputDelta,const CTensor<type_t> &cTensor_Delta,const std::vector<CTensor<type_t> > &cTensor_Kernel,const std::vector<type_t> &bias)
{
 int32_t padding_x=0;//дополнение нулями
 int32_t padding_y=0;//дополнение нулями
 int32_t step_x=1;//шаг свёртки
 int32_t step_y=1;//шаг свёртки

 //вычисляем размеры выходного тензора
 int32_t kernel_amount=cTensor_Kernel.size();
 if (kernel_amount==0) throw("Для обратной свёртки требуется хотя бы одно ядро свёртки");
 if (kernel_amount!=bias.size()) throw("Для обратной свёртки требуется чтобы количество ядер и смещений совпадало");

 int32_t kernel_x=cTensor_Kernel[0].Size_X;
 int32_t kernel_y=cTensor_Kernel[0].Size_Y;
 int32_t kernel_z=cTensor_Kernel[0].Size_Z;

 int32_t input_y=cTensor_Delta.Size_Y;
 int32_t input_x=cTensor_Delta.Size_X;
 int32_t input_z=cTensor_Delta.Size_Z;

 //обратная свёртка делается с ядрами, повёрнутыми на 180
 int32_t output_x=step_x*(input_x-1)+kernel_x-2*padding_x;
 int32_t output_y=step_y*(input_y-1)+kernel_y-2*padding_y;
 int32_t output_z=kernel_z;

 padding_x=kernel_x-1-padding_x;
 padding_y=kernel_y-1-padding_y;

 if (cTensor_OutputDelta.Size_X!=output_x || cTensor_OutputDelta.Size_Y!=output_y || cTensor_OutputDelta.Size_Z!=output_z) throw("Ошибочная размерность выходного тензора для обратной свёртки");
 if (input_z!=kernel_amount) throw("Для обратной свёртки требуется чтобы количество фильтров и глубина входного тензора совпадали");

 STensorKernel<type_t> sTensorKernel_OutputDelta(cTensor_OutputDelta);//выходной тензор
 STensorKernel<type_t> sTensorKernel_Delta(cTensor_Delta);//входной тензор

 cTensor_Delta.CopyToDevice();

 //копируем на видеокарту указатели и смещения
 CCUDADeviceVector<type_t> bias_array(bias.size());
 bias_array.copy_host_to_device(&bias[0],bias.size());

 CCUDADeviceVector<type_t *> kernel_item_ptr_array(cTensor_Kernel.size());
 std::vector<type_t *> item_ptr(cTensor_Kernel.size());
 for(size_t n=0;n<cTensor_Kernel.size();n++)
 {
  cTensor_Kernel[n].CopyToDevice();
  item_ptr[n]=cTensor_Kernel[n].GetDeviceVector().get();
 }
 kernel_item_ptr_array.copy_host_to_device(&item_ptr[0],item_ptr.size());
 //выполняем свёртку
 STensorKernel<type_t> sTensorKernel_Kernel(cTensor_Kernel[0]);

 dim3 thread(CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE,CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE);

 size_t block_x=output_x/thread.x;
 if (output_x%thread.x) block_x++;
 size_t block_y=output_y/thread.y;
 if (output_y%thread.y) block_y++;
 size_t block_z=output_z;

 dim3 blocks(block_x,block_y,block_z);
 if (blocks.x==0) blocks.x=1;
 if (blocks.y==0) blocks.y=1;
 if (blocks.z==0) blocks.z=1;
 CUDABackwardConvolutionFunction<type_t><<<blocks,thread>>>(sTensorKernel_OutputDelta,sTensorKernel_Delta,sTensorKernel_Kernel,kernel_item_ptr_array.get(),cTensor_Kernel.size(),bias_array.get(),padding_x,padding_y,step_x,step_y);
 HANDLE_ERROR(cudaGetLastError());
 HANDLE_ERROR(cudaDeviceSynchronize());

/*
 dim3 thread(output_z,1,1);
 dim3 blocks(output_x,output_y,1);
 if (blocks.x==0) blocks.x=1;
 if (blocks.y==0) blocks.y=1;
 if (blocks.z==0) blocks.z=1;
 CUDABackwardConvolutionFunction<type_t><<<blocks,thread>>>(sTensorKernel_OutputDelta,sTensorKernel_Delta,sTensorKernel_Kernel,kernel_item_ptr_array.get(),cTensor_Kernel.size(),bias_array.get(),padding_x,padding_y,step_x,step_y);
 HANDLE_ERROR(cudaGetLastError());
 HANDLE_ERROR(cudaDeviceSynchronize());
*/
 cTensor_OutputDelta.SetDeviceOnChange();

/*
 for(int32_t y=0;y<output_y;y++)
 {
  for(int32_t x=0;x<output_x;x++)
  {
   for(int32_t z=0;z<output_z;z++)
   {
    type_t summ=0;
    //идём по всем весовым коэффициентам фильтров
    for(int32_t ky=0;ky<kernel_y;ky++)
    {
     for(int32_t kx=0;kx<kernel_x;kx++)
     {
      int32_t y0=y+ky-padding_y;
      int32_t x0=x+kx-padding_x;
      //игнорируем выходящие за границы элементы
      if (y0<0 || y0>=input_y) continue;
      if (x0<0 || x0>=input_x) continue;
      //суммируем по всем фильтрам
      for(int32_t f=0;f<kernel_amount;f++)
      {
       type_t k=cTensor_Kernel[f].GetElement(z,kernel_y-1-ky,kernel_x-1-kx);
       type_t d=cTensor_Delta.GetElement(f,y0,x0);
       summ+=k*d;
      }
     }
    }
    cTensor_OutputDelta.SetElement(z,y,x,summ);
   }
  }
 }
 cTensor_OutputDelta.SetHostOnChange();
*/

 /*
 for(int32_t y=0;y<output_y;y++)
 {
  for(int32_t x=0;x<output_x;x++)
  {
   for(int32_t z=0;z<output_z;z++)
   {
    const type_t *delta_ptr=cTensor_Delta.GetColumnPtr(0,0);
    type_t *output_ptr=cTensor_OutputDelta.GetColumnPtr(0,0)+z*output_x*output_y+y*output_x+x;
    size_t kernel_depth_offset=z*kernel_x*kernel_y;
    type_t sum=0;//сумма для градиента
    //идём по всем весовым коэффициентам фильтров
    for(size_t ky=0;ky<kernel_y;ky++)
    {
     int32_t y0=static_cast<int32_t>(y*step_y+ky);//TODO: возможно, ошибочно умножать на шаг
     y0-=static_cast<int32_t>(padding_y);
     if (y0<0 || y0>=input_y) continue;
     for(size_t kx=0;kx<kernel_x;kx++)
     {
      int32_t x0=static_cast<int32_t>(x*step_x+kx);//TODO: возможно, ошибочно умножать на шаг
      x0-=static_cast<int32_t>(padding_x);
      //игнорируем выходящие за границы элементы
      if (x0<0 || x0>=input_x) continue;
      //суммируем по всем ядрам
      size_t offset_k_ptr=(kernel_y-1-ky)*kernel_x+(kernel_x-1-kx)+kernel_depth_offset;
      size_t offset_d_ptr=y0*input_x+x0;
      for(size_t k=0;k<kernel_amount;k++)
      {
       sum+=bias[k];//TODO: надо выяснить, как прибавлять смещения
       const type_t *d_ptr=delta_ptr+k*input_x*input_y+offset_d_ptr;
       const type_t *kernel_ptr=cTensor_Kernel[k].GetColumnPtr(0,0);
       const type_t *k_ptr=kernel_ptr+offset_k_ptr;
       sum+=(*k_ptr)*(*d_ptr);//добавляем произведение повёрнутых фильтров на дельты
      }
     }
    }
    *output_ptr=sum;//записываем результат в тензор градиента
   }
  }
 }

 cTensor_OutputDelta.SetHostOnChange();
 */

}

//----------------------------------------------------------------------------------------------------
//функция CUDA для выполнения поправок
//----------------------------------------------------------------------------------------------------
template<class type_t>
__global__ void CUDADeltaWeightAndBiasFunction(STensorKernel<type_t> tensor_d_kernel,type_t** d_kernel_item_ptr_array,type_t* d_bias_ptr,STensorKernel<type_t> tensor_image,STensorKernel<type_t> tensor_delta,int32_t padding_x,int32_t padding_y,int32_t step_x,int32_t step_y,size_t dkernel_amount)
{
 size_t blockCol=blockIdx.x;
 size_t blockRow=blockIdx.y;
 size_t kz=blockIdx.z/dkernel_amount;
 //координаты элементов блока в выходном тензоре
 size_t kx=threadIdx.x;
 size_t ky=threadIdx.y;

 kx+=blockCol*CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE;
 ky+=blockRow*CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE;

 int32_t d_kernel_x=tensor_d_kernel.GetSizeX();
 int32_t d_kernel_y=tensor_d_kernel.GetSizeY();

 if (kx>=d_kernel_x) return;
 if (ky>=d_kernel_y) return;

 int32_t f=blockIdx.z%dkernel_amount;
/*
 int32_t kx=blockIdx.x;
 int32_t ky=blockIdx.y;
 int32_t kz=blockIdx.z;

 int32_t f=threadIdx.x;
*/
 int32_t image_x=tensor_image.GetSizeX();
 int32_t image_y=tensor_image.GetSizeY();

 int32_t delta_x=tensor_delta.GetSizeX();
 int32_t delta_y=tensor_delta.GetSizeY();
 int32_t delta_z=tensor_delta.GetSizeZ();

 //настроим ядро
 tensor_d_kernel.TensorData_Ptr=d_kernel_item_ptr_array[f];

 for(int32_t y=0;y<delta_y;y++)
 {
  for(int32_t x=0;x<delta_x;x++)
  {
   type_t delta=tensor_delta.GetElement(f,y,x);//запоминаем значение градиента
   int32_t i0=ky+y*step_y-padding_y;//TODO: возможно, ошибочно умножать на шаг
   int32_t j0=kx+x*step_x-padding_x;//TODO: возможно, ошибочно умножать на шаг
   if (i0>=0 && i0<image_y && j0>=0 && j0<image_x)//игнорируем выходящие за границы элементы
   {
    //наращиваем градиент фильтра
    type_t dk=tensor_d_kernel.GetElement(kz,ky,kx);
    dk+=delta*tensor_image.GetElement(kz,i0,j0);
    tensor_d_kernel.SetElement(kz,ky,kx,dk);
   }
   if (kx==0 && ky==0 && kz==0) d_bias_ptr[f]+=delta;
  }
 }
}

//----------------------------------------------------------------------------------------------------
/*!создание поправок весов и смещений
*/
//----------------------------------------------------------------------------------------------------
template<class type_t>
void CTensorConv<type_t>::CreateDeltaWeightAndBias(std::vector<CTensor<type_t> > &cTensor_dKernel,std::vector<type_t> &dbias,const CTensor<type_t> &cTensor_Image,const CTensor<type_t> &cTensor_Delta)
{
 int32_t padding_x=0;//дополнение нулями
 int32_t padding_y=0;//дополнение нулями
 int32_t step_x=1;//шаг свёртки
 int32_t step_y=1;//шаг свёртки

 int32_t dkernel_amount=cTensor_dKernel.size();
 if (dkernel_amount==0) throw("Для создания поправок весов и смещений требуется не пустой вектор поправок к ядрам");
 if (dbias.size()!=dkernel_amount) throw("Для создания поправок весов и смещений требуется чтобы количество поправок фильтров и поправок сдвигов совпадало");

 //TODO: считать все тензоры с видеокарты

 int32_t image_x=cTensor_Image.Size_X;
 int32_t image_y=cTensor_Image.Size_Y;
 int32_t image_z=cTensor_Image.Size_Z;

 int32_t delta_x=cTensor_Delta.Size_X;
 int32_t delta_y=cTensor_Delta.Size_Y;
 int32_t delta_z=cTensor_Delta.Size_Z;

 int32_t dkernel_x=image_x-delta_x+1;
 int32_t dkernel_y=image_y-delta_y+1;
 int32_t dkernel_z=image_z;

 if (dkernel_x!=cTensor_dKernel[0].Size_X || dkernel_y!=cTensor_dKernel[0].Size_Y || dkernel_z!=cTensor_dKernel[0].Size_Z) throw("Неверные размеры тензора поправок к ядрам для обновления весов и смещений");

 cTensor_Delta.CopyToDevice();
 cTensor_Image.CopyToDevice();

 STensorKernel<type_t> sTensorKernel_Delta(cTensor_Delta);
 STensorKernel<type_t> sTensorKernel_Image(cTensor_Image);

 CCUDADeviceVector<type_t> d_bias_array(dbias.size());
 d_bias_array.copy_host_to_device(&dbias[0],dbias.size());

 CCUDADeviceVector<type_t *> d_kernel_item_ptr_array(cTensor_dKernel.size());
 std::vector<type_t *> item_ptr(cTensor_dKernel.size());
 for(size_t n=0;n<cTensor_dKernel.size();n++)
 {
  cTensor_dKernel[n].CopyToDevice();
  item_ptr[n]=cTensor_dKernel[n].GetDeviceVector().get();
 }
 d_kernel_item_ptr_array.copy_host_to_device(&item_ptr[0],item_ptr.size());
 //выполняем свёртку
 STensorKernel<type_t> sTensorKernel_dKernel(cTensor_dKernel[0]);

 dim3 thread(CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE,CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE,1);

 size_t block_x=dkernel_x/thread.x;
 if (dkernel_x%thread.x) block_x++;
 size_t block_y=dkernel_y/thread.y;
 if (dkernel_y%thread.y) block_y++;
 size_t block_z=dkernel_amount+dkernel_z*dkernel_amount;

 dim3 blocks(block_x,block_y,block_z);
 if (blocks.x==0) blocks.x=1;
 if (blocks.y==0) blocks.y=1;
 if (blocks.z==0) blocks.z=1;
 CUDADeltaWeightAndBiasFunction<type_t><<<blocks,thread>>>(sTensorKernel_dKernel,d_kernel_item_ptr_array.get(),d_bias_array.get(),sTensorKernel_Image,sTensorKernel_Delta,padding_x,padding_y,step_x,step_y,dkernel_amount);
 HANDLE_ERROR(cudaGetLastError());
 HANDLE_ERROR(cudaDeviceSynchronize());

/*
 dim3 thread(dkernel_amount,1,1);
 dim3 blocks(dkernel_x,dkernel_y,dkernel_z);
 if (blocks.x==0) blocks.x=1;
 if (blocks.y==0) blocks.y=1;
 if (blocks.z==0) blocks.z=1;
 CUDADeltaWeightAndBiasFunction<type_t><<<blocks,thread>>>(sTensorKernel_dKernel,d_kernel_item_ptr_array.get(),d_bias_array.get(),sTensorKernel_Image,sTensorKernel_Delta,padding_x,padding_y,step_x,step_y);
 HANDLE_ERROR(cudaGetLastError());
 HANDLE_ERROR(cudaDeviceSynchronize());*/

 for(size_t n=0;n<cTensor_dKernel.size();n++)
 {
  cTensor_dKernel[n].SetDeviceOnChange();
 }
 d_bias_array.copy_device_to_host(&dbias[0],dbias.size());

/*
 for(int32_t f=0;f<dkernel_amount;f++)
 {
  cTensor_dKernel[f].SetHostOnChange();

  for(int32_t y=0;y<delta_y;y++)
  {
   for(int32_t x=0;x<delta_x;x++)
   {
    type_t delta=cTensor_Delta.GetElement(f,y,x);//запоминаем значение градиента
    for(int32_t i=0;i<dkernel_y;i++)
    {
     for(int32_t j=0;j<dkernel_x;j++)
     {
      int32_t i0=i+y*step_y-padding_y;//TODO: возможно, ошибочно умножать на шаг
      int32_t j0=j+x*step_x-padding_x;//TODO: возможно, ошибочно умножать на шаг
      //игнорируем выходящие за границы элементы
      if (i0<0 || i0>=image_y || j0<0 || j0>=image_x) continue;
      //наращиваем градиент фильтра
      for(int32_t c=0;c<dkernel_z;c++)
      {
       type_t dk=cTensor_dKernel[f].GetElement(c,i,j);
       dk+=delta*cTensor_Image.GetElement(c,i0,j0);
       cTensor_dKernel[f].SetElement(c,i,j,dk);
      }
     }
    }
    dbias[f]+=delta;
   }
  }
 }*/

}

//----------------------------------------------------------------------------------------------------
//функция CUDA для выполнения поправок для обратной свёртки
//----------------------------------------------------------------------------------------------------
template<class type_t>
__global__ void CUDABackDeltaWeightAndBiasFunction(STensorKernel<type_t> tensor_d_kernel,type_t** d_kernel_item_ptr_array,type_t* d_bias_ptr,STensorKernel<type_t> tensor_image,STensorKernel<type_t> tensor_delta,int32_t padding_x,int32_t padding_y,int32_t step_x,int32_t step_y,int32_t dkernel_amount)
{
 size_t blockCol=blockIdx.x;
 size_t blockRow=blockIdx.y;
 size_t kz=blockIdx.z/dkernel_amount;
 //координаты элементов блока в выходном тензоре
 size_t kx=threadIdx.x;
 size_t ky=threadIdx.y;

 kx+=blockCol*CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE;
 ky+=blockRow*CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE;

 int32_t d_kernel_x=tensor_d_kernel.GetSizeX();
 int32_t d_kernel_y=tensor_d_kernel.GetSizeY();

 if (kx>=d_kernel_x) return;
 if (ky>=d_kernel_y) return;

 int32_t f=blockIdx.z%dkernel_amount;

/*
 int32_t kx=blockIdx.x;
 int32_t ky=blockIdx.y;
 int32_t kz=blockIdx.z;

 int32_t f=threadIdx.x;
*/
 int32_t image_x=tensor_image.GetSizeX();
 int32_t image_y=tensor_image.GetSizeY();

 int32_t delta_x=tensor_delta.GetSizeX();
 int32_t delta_y=tensor_delta.GetSizeY();
 int32_t delta_z=tensor_delta.GetSizeZ();

 //настроим ядро
 tensor_d_kernel.TensorData_Ptr=d_kernel_item_ptr_array[f];

 for(int32_t y=0;y<image_y;y++)
 {
  for(int32_t x=0;x<image_x;x++)
  {
   type_t image=tensor_image.GetElement(f,y,x);//запоминаем значение градиента
   int32_t i0=ky+y*step_y-padding_y;//TODO: возможно, ошибочно умножать на шаг
   int32_t j0=kx+x*step_x-padding_x;//TODO: возможно, ошибочно умножать на шаг
   if (i0>=0 && i0<delta_y && j0>=0 && j0<delta_x)//игнорируем выходящие за границы элементы
   {
    //наращиваем градиент фильтра
    type_t dk=tensor_d_kernel.GetElement(kz,ky,kx);
    dk+=image*tensor_delta.GetElement(kz,i0,j0);
    tensor_d_kernel.SetElement(kz,ky,kx,dk);
   }
  }
 }

}

//----------------------------------------------------------------------------------------------------
/*!создание поправок весов и смещений для обратной свёртки
*/
//----------------------------------------------------------------------------------------------------
template<class type_t>
void CTensorConv<type_t>::CreateBackDeltaWeightAndBias(std::vector<CTensor<type_t> > &cTensor_dKernel,std::vector<type_t> &dbias,const CTensor<type_t> &cTensor_Image,const CTensor<type_t> &cTensor_Delta)
{
 int32_t padding_x=0;//дополнение нулями
 int32_t padding_y=0;//дополнение нулями
 int32_t step_x=1;//шаг свёртки
 int32_t step_y=1;//шаг свёртки

 int32_t dkernel_amount=cTensor_dKernel.size();
 if (dkernel_amount==0) throw("Для создания поправок весов и смещений требуется не пустой вектор поправок к ядрам");
 if (dbias.size()!=dkernel_amount) throw("Для создания поправок весов и смещений требуется чтобы количество поправок фильтров и поправок сдвигов совпадало");

 int32_t delta_x=cTensor_Delta.Size_X;
 int32_t delta_y=cTensor_Delta.Size_Y;
 int32_t delta_z=cTensor_Delta.Size_Z;

 int32_t image_x=cTensor_Image.Size_X;
 int32_t image_y=cTensor_Image.Size_Y;
 int32_t image_z=cTensor_Image.Size_Z;

 int32_t dkernel_x=delta_x-image_x+1;
 int32_t dkernel_y=delta_y-image_y+1;
 int32_t dkernel_z=delta_z;

 if (dkernel_x!=cTensor_dKernel[0].Size_X || dkernel_y!=cTensor_dKernel[0].Size_Y || dkernel_z!=cTensor_dKernel[0].Size_Z) throw("Неверные размеры тензора поправок к ядрам для обновления весов и смещений");

 cTensor_Delta.CopyToDevice();
 cTensor_Image.CopyToDevice();

 STensorKernel<type_t> sTensorKernel_Delta(cTensor_Delta);
 STensorKernel<type_t> sTensorKernel_Image(cTensor_Image);

 CCUDADeviceVector<type_t> d_bias_array(dbias.size());
 d_bias_array.copy_host_to_device(&dbias[0],dbias.size());

 CCUDADeviceVector<type_t *> d_kernel_item_ptr_array(cTensor_dKernel.size());
 std::vector<type_t *> item_ptr(cTensor_dKernel.size());
 for(size_t n=0;n<cTensor_dKernel.size();n++)
 {
  cTensor_dKernel[n].CopyToDevice();
  item_ptr[n]=cTensor_dKernel[n].GetDeviceVector().get();
 }
 d_kernel_item_ptr_array.copy_host_to_device(&item_ptr[0],item_ptr.size());
 //выполняем свёртку
 STensorKernel<type_t> sTensorKernel_dKernel(cTensor_dKernel[0]);

 dim3 thread(CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE,CTensorMath<type_t>::TENSOR_OPERATION_BLOCK_SIZE,1);

 size_t block_x=dkernel_x/thread.x;
 if (dkernel_x%thread.x) block_x++;
 size_t block_y=dkernel_y/thread.y;
 if (dkernel_y%thread.y) block_y++;
 size_t block_z=dkernel_amount+dkernel_z*dkernel_amount;

 dim3 blocks(block_x,block_y,block_z);
 if (blocks.x==0) blocks.x=1;
 if (blocks.y==0) blocks.y=1;
 if (blocks.z==0) blocks.z=1;
 CUDABackDeltaWeightAndBiasFunction<type_t><<<blocks,thread>>>(sTensorKernel_dKernel,d_kernel_item_ptr_array.get(),d_bias_array.get(),sTensorKernel_Image,sTensorKernel_Delta,padding_x,padding_y,step_x,step_y,dkernel_amount);
 HANDLE_ERROR(cudaGetLastError());
 HANDLE_ERROR(cudaDeviceSynchronize());

 for(size_t n=0;n<cTensor_dKernel.size();n++)
 {
  cTensor_dKernel[n].SetDeviceOnChange();
 }
 d_bias_array.copy_device_to_host(&dbias[0],dbias.size());
}

VDG
18.06.2023 05:58
#25663434
размерность полной матрицы вы в него не вставите никогда для матриц превышающих по размеру число потоков

Не знаю как там в CUDA, пишу прямо в шейдерах. Но зачем пытаться запихнуть всю матрицу в блок? Каждая ячейка считает своё скалярное произведение и пишет результат в память выходной матрицы, ориентируясь на свои координаты в ней.

Размер блока надо подбирать кратным размерам выходной матрицы, чтобы данные, к которым обращаются ячейки, были по возможности одни и те же, и оставались в кеше без его переполнения и повторного считывания из основной памяти карты.
1. da-nie
  18.06.2023 05:58
  #25663610
  Но зачем пытаться запихнуть всю матрицу в блок?
  
  Имелось в виду, что использовать эти размерности потоков для индексации элементов матрицы нельзя. То есть, отдать потоку элементы всей матрицы нельзя (матрицы можно будет использовать только очень маленькие).
  
  и оставались в кеше
  
  Там кэш — та самая разделяемая память и он с ручным управлением. То есть, что туда положите, то и в кэше. Так вот, разбиение на блоки 16x16 как раз и призвано использовать кэщ 16x16 элементов.

Sadler
18.06.2023 05:58
#25663266
Или по сетям GAN как считать ошибку выходного слоя дискриминатора.

Попробуйте Wasserstein GAN.Оно намного стабильнее учится, нежели generic GAN, и метод расчёта loss-ов явно указан в статье, не надо ничего выискивать.
1. da-nie
  18.06.2023 05:58
  #25663280
  Спасибо! Посмотрю.

iskateli
18.06.2023 05:58
#25663118
+1
а разве только с параметрами работает нейросеть?

Параметрическая адаптация означает, что подстраиваются числовые параметры модели, минимизируются различные ошибки и пр.

Структурная адаптация - означает, что модель не работает с символами (числами), а работает по принципу адаптивного пространственно-временного фильтра для входящих потоков данных, просто перенаправляя потоки в правильном направлении.

grigorym
18.06.2023 05:58
#25663690
+2
Вашей активационной функцией f(x) = x пренебречь можно, согласно формуле, не только при положительных, но и при отрицательных значениях x. А значит выход всей вашей нейронной сети является линейной комбинацией входных параметров и грош ей цена, ничего делать она не сможет.