Comments 28
Это очень здорово, но гораздо интереснее было бы рассказать, как именно происходит обучение пакетами на GPU.
1) Весь пакет прогоняется на GPU параллельно? Обычно, столько памяти нет.
2) На CPU сделан цикл по элементам пакета. Один элемент копируется на GPU и производится обучение на нём с сумированием весов. В этом случае непонятно, как делать нормализацию?
3) На GPU копируется пакет и один слой сети. Выполняется проход по одному слою сети для всех элементов пакета. Далее, пакет выбрасывается и выход слоя является новым пакетом для следующего слоя. Памяти нужно меньше, чем в пункте 1, но все промежуточные результаты придётся копировать на CPU (они нужны для обратного прохода). А такое копирование очень медленное.
Судя по минусаторам, все они, конечно, знают, как же сделано на практике в пакетах типа keras и tensorflow? Не так ли? ;) Но при этом они не знают по какой математике считается нейрон, раз аплодируют этой абсолютно примитивной статье?
интереснее было бы рассказать, как именно происходит обучение
поэтому такие этапы как обучение и т.д. в ней не затрагиваются
интересно это было написано до вашего комментария? но в любом случае напрасно опускать такой важный этап как обучение, раз "Цель данной публикации – комплексное рассмотрение "
Нейрон - это взвешивающий сумматор с (опционально) нелинейным выходом. Хорошо, если функция активации дифференцируема на R. Тогда его можно учить методом обратного распространения ошибки. Если нет - таким способом учить его невозможно (только Больцманом - закреплением случайных флуктуаций весов по локально-лучшему результату)
Тут автору каждой сотой статьи про то как работают нейронные сети ачивку выдают, или в чем смысл их постоянно постить?)
Каждый автор считает своим долгом донести внезапно полученные знания о нейросетях (часто очень примитивные). :) Жаль только, что в математике и принципах обучения они редко что-то понимают, ибо популярные пакеты всё это тщательно скрывают. В результате по тем же свёрточным сетям хрен найдёшь как их обучать. Или по сетям GAN как считать ошибку выходного слоя дискриминатора. Или как делать нормализацию на пальцах. Или как же обучать пакетом на GPU. Или как сделать быструю свёртку, обратную свёртку и коррекцию коэффициентов на GPU. А в остальном да, все специалисты. :)
А здесь статья с очень громким названием "Что под капотом у нейронной сети. Нейросеть c точки зрения математики и программирования". То есть, ожидается как минимум вся подноготная. Но… имеем то, что имеем. Мой вон вопрос-предложение выше вызывает у многих баттхерт. А я всего лишь спросил ту информацию, которую я не сумел пока ещё нигде найти.
Информацию в смысле как на GPU много данных обсчитывается? Если да - то ответ прост - батчами, т.е. не все примеры сразу загоняются в память, а столько, сколько влазит, после каждого прогона батча обновляются веса. Да, при этом градиент считается не по всему обучающему множеству сразу, но зато в память влазит. Мне кажется во всех современных нейросетевых библиотеках есть параметр batch_size, который надо подгонять под то чтобы побольше памяти за раз использовать, но не поймать Out Of Memory. Как то так.
Нет, ответ не прост. Даже для достаточно примитивных сетей требуется весьма много памяти, и пакет у вас будет очень маленьким (и вообще выродится в один образ). Потом, от размера пакета будет зависеть нормализация.
Вот тут я делал на С++ такую сеть с самодельной работой с GPU.
Так вот, в мою карту 1060 с 3 ГБ поместилось за раз не более 100 входных изображений из 60000. Для более сложной сети вообще помещается один образ и не более. И как нормализацию делать в таком случае? Вот именно.
Ну 100 - это уже хорошо, на то они и мини-батчи. А один пример за раз - тоже не новость, называется "online learning". Вот тут сравнение всех трех вариантов с примерами и т.п.: Full batch, mini-batch, and online learning
Так 100 — это для очень-очень простой сети. Когда я 160x120 генерировать пытался на свёрточной сети, там вообще штук 5 образов влезало за раз. И всю память жрут полносвязаные слои (матрицы весов быстро нарастают). Сами ядра свёрток мало жрут.
Поэтому мне вот кажется, вероятно, послойно там всё обрабатывается внутри с выбросом использованных слоёв.
Есть ещё вопросы к быстрой свёртке. Вот есть алгоритм свёртки через умножение. Я его сделал влоб сперва. Какой наивный! :) Одна (!) матрица такой свёртки для картинки 300x300 весила больше 8 ГБ. Ладно, переделал с генерацией матрицы на лету и умножением матрицы с shared-памятью. Заработало. Но скорость упала в 100 раз по сравнению с простой распараллеленой свёрткой без shared-памяти на GPU — оно и понятно, очень много лишних умножений. Мда.
Кстати! Вот смотрите, есть у вас CUDA. Там у вас есть трёхмерные блоки и трёхмерные потоки внутри блоков. Смотрите какая штука вас ожидает.
Быстрое матричное умножение (а оно потребуется для полносвязных сетей) требует задать в потоках блок x=16, y=16,z=1 (число всех потоков внутри блока сильно ограничено! 512, 1024 и, может, уже чуть больше на новых картах. То есть, размерность полной матрицы вы в него не вставите никогда для матриц превышающих по размеру число потоков). Осталась размерность по Z, но что там поместится? 1024/(256)=4. Вот максимум размерности по Z. Вообще не используем её в этом случае. Как видите, мы уже заняли размерность потока. Теперь берём блоки. В x и y мы помещаем номер блока 16x16 (который обрабатывают потоки) по x и по y внутри матрицы результата умножения. А что в z? Номер батча? Ну уж нет. У нас не матрицы, а 3D-тензоры в общем случае, поэтому там z для тензора.
То есть, мы уже заняли всю параллельную работу на видяхе на одну только матрицу. :)
Ну если так микроскопом гвозди забивать - конечно никаких ресурсов не хватит ?Откуда матрицы свертки на 8 Gb? Обычно они по сотне байтов занимают (зависит от размера, конечно). Я на MNIST первую сверточную сеть запускал LeNet или AlexNet - не помню, но она обучилась нормально на двухядерном процессоре с 4Gb памяти, без всяких GPU. Посмотрели бы как это сделано в том же TensorFlow - он на C++ написан, а на питоне только биндинги для удобного взаимодействия (а Keras - это ещё немного абстракций на питоне для того же TensorFlow, не более). PyTorch - так же, как и numpy и т.д.
Вот откуда:
Или так нагляднее:
Имея быстрое умножение матриц на GPU (на shared-памяти) очень хочется его использовать для ускорения свётки (так как shared-память я пока использовать для произвольных параметров свёртки не придумал как).
И вот есть статья.
Посмотрели бы как это сделано в том же TensorFlow
Это не так-то просто будет понять.
ого, вот это вы мощно - это ж сколько движений для создания матриц и перемножения!
а вы не пробовали
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] ...)
это ж сколько движений для создания матриц и перемножения!
Это только вариант на аналогах всяких матриц Теплица. Оказался неудачным.
а вы не пробовали
Я пока не очень понял, что вы имеете в виду. В исходном варианте и g и f не вектора, а матрицы. Их напрямую так не перемножишь.
Я сейчас пробую понять, что сделано тут.
насколько я понял, там делают предрасчет сумм окна свертки (ибо это можно сделать одним проходом, очень эффективно), а потом уже умножают на коэфициенты
Меня смущает вот что:
Исходя из того, что K = srcC kernelY kernelX / group, эффективность метода особенно низка для входных свёрточных слоев. А для depthwise convolution матричный метод вообще проигрывает тривиальной реализации.
Дело ещё в том, что кроме прямой и обратной свётки (и перевёрнутой) есть ещё расчёт поправок коэффициентов ядер. Этот этап тоже надо думать как ускорить.
Сейчас сделал вот так:
Работает быстро. Но памяти тоже надо дофига.
Ну и какая-же это оптимизация, если вместо маленькой матрицы, которую можно вообще в константы шейдера задать, у вас 8Gb чисел ...
Задать свертку шейдером и запустить на просчет на каждом участке видеокарты - оно как раз распараллелится на кол-во CUDA-ядер, и вся свертка в два присеста посчитается, а памяти надо будет порядка размера обрабатываемого изображения.
Как я уже раза два написал, это вариант через умножение матриц.
Задать свертку шейдером
Не использовал шейдеры, но, полагаю, максимального быстродействия вы не получите, если просто сделаете свёртку шейдером.
#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());
}Вообще, согласно статье свёртка и обратная свёртка через умножение матриц делается так:
Я сейчас так и сделал. Работает быстро. Но памяти тоже надо дофига.
размерность полной матрицы вы в него не вставите никогда для матриц превышающих по размеру число потоков
Не знаю как там в CUDA, пишу прямо в шейдерах. Но зачем пытаться запихнуть всю матрицу в блок? Каждая ячейка считает своё скалярное произведение и пишет результат в память выходной матрицы, ориентируясь на свои координаты в ней.
Размер блока надо подбирать кратным размерам выходной матрицы, чтобы данные, к которым обращаются ячейки, были по возможности одни и те же, и оставались в кеше без его переполнения и повторного считывания из основной памяти карты.
Но зачем пытаться запихнуть всю матрицу в блок?
Имелось в виду, что использовать эти размерности потоков для индексации элементов матрицы нельзя. То есть, отдать потоку элементы всей матрицы нельзя (матрицы можно будет использовать только очень маленькие).
и оставались в кеше
Там кэш — та самая разделяемая память и он с ручным управлением. То есть, что туда положите, то и в кэше. Так вот, разбиение на блоки 16x16 как раз и призвано использовать кэщ 16x16 элементов.
Или по сетям GAN как считать ошибку выходного слоя дискриминатора.
Попробуйте Wasserstein GAN.Оно намного стабильнее учится, нежели generic GAN, и метод расчёта loss-ов явно указан в статье, не надо ничего выискивать.
а разве только с параметрами работает нейросеть?
Параметрическая адаптация означает, что подстраиваются числовые параметры модели, минимизируются различные ошибки и пр.
Структурная адаптация - означает, что модель не работает с символами (числами), а работает по принципу адаптивного пространственно-временного фильтра для входящих потоков данных, просто перенаправляя потоки в правильном направлении.
Вашей активационной функцией f(x) = x пренебречь можно, согласно формуле, не только при положительных, но и при отрицательных значениях x. А значит выход всей вашей нейронной сети является линейной комбинацией входных параметров и грош ей цена, ничего делать она не сможет.
Что под капотом у нейронной сети. Нейросеть c точки зрения математики и программирования