05. Отчет по первому опыту работы с технологией CUDA в реализации искусственной нейронной сети

Введение

Начав изучать C/C++, я неизбежно пришел к теме многопоточного и параллельного программирования и высокопроизводительных вычислений. Почитав Рихтера [1] и поупражнявшись с многопоточностью в Windows, я начал искать способы организации многопоточности попроще и поэффективнее.

После изучения технологии OpenMP, которая привела меня в восторг, PPL, MPI, я подошел к технологиям GPGPU (использование графического видеоускорителя для вычислений общего назначения). Почитав отзывы, я решил попробовать ее в деле.

На настоящий момент (2014 год) имелись следующие технологии GPGPU:

• • CUDA (Compute Unified Device Architecture) от Nvidia.

  • AMD FireStream от (как нетрудно догадаться) AMD

  • C++ AMP (Accelerated Massive Parallelism) от Microsoft.

  • OpenCL (Open Computing Language) - свободный фреймворк, не привязывающий к определенной видеокарте.

  • OpenACC - стандарт, аналогичный OpenMP, но предназначенный для использования с GPU.

Технология AMP очень хороша, по крайней мере, на сайте msdn Microsoft и в статьях, [1]. Использовать ее мне мешает только отсутствие MS Visual Studio 2012 (использую в работе MS VS 2010). OpenACC еще не мог достать. OpenCL, судя по литературе - неэффективный и плохо развиваемый. Оставались только CUDA и FireStream. Из них я выбрал CUDA чисто по личным предпочтениям и еще, потому что в наличии имеется достаточное количество литературы.

Тестовая задача

В качестве тестовой задачи было решено выбрать реализацию перцептрона. Перцептрон - одна из старейших компьютерных моделей мозга, предложенная Фрэнком Розенбаттом аж в лохматом 1957 году. Описывать алгоритм я не буду, все равно Вы уже посмотрели его в Википедии. Вычисления в перцептроне (и вообще в искуственной нейронной сети (ИНС) состоят в перемножении элементов двух векторов и их сложения. Алгоритм довольно легко подвергается распараллеливанию.

Еще одной причиной выбора такой тестовой задачи является увлечение автора технлогиями ИИ, отсутствие реализации перцептрона в примерах к CUDA SDK и литературе, а также перспективностью нейросетевых технологий вообще.

Цели эксперимента.

Мне хотелось сравнить работу контейнеров stl в сравнении с массивами. Выяснить реальные потери производительности при использовании контейнеров. Контейнеры хорошы, когда требуется частое добавление-удаление элементов. Но какая производительность приносится в жертву удобству работы с контейнерами? Это необходимо выяснить.

• 1. Отработать технологию CUDA (как частный случай GPGPU) на простом примере перцептрона.

  2. Отработать технологию ИНС на примере перцептрона.

  3. Сравнить производительность последовательных реализаций на основе контейнеров stl и массивов.

  4. Сравнить производительность разработки с использованием массивов/контейнеров.

  5. Сравнить производительность разработки с использованием C/C++ CPU и C CUDA

  6. Создать задел кода на будущее для реализации полноценных ИНС.

План работы

Необходимо написать несколько реализаций вычисления перцептрона. Алгоритм обучения предлагается пока не реализовывать, так как считаю, что оценить потенциал GPGPU можно уже по функции вычисления перцептрона.

Реализовывать следующими технологиями.

• • На основе std::Vector в последовательной реализации.

  • На основе std::Vector в параллельной реализации с использованием PPL.

  • На основе обычных массивов в последовательной реализации.

  • На основе обычных массивов с использованием CUDA.

Технические решения

Выбор оборудования

Целью данной работы стояло поработать с технологией CUDA. Максимальная производительность не требовалось, но и оборудование совсем уж начального уровня брать не хотелось. В то же время цена должна быть минимальной, чтобы не жалко было выбросить карту вследствии естественного морального износа. Сравнив процессоры в статье на Википедии, в качестве оборудования была выбрана видеокарта на основе GeForce GT 640. Что интересно, модель GT 640 GDDR обладала более высокими характеристиками памяти. Однако для нее не указывалась поддежка CUDA, поэтому приобретать ее я поостерегся и возможно, зря.

Инструменты и программное обеспечение

Для разработки с использованием CUDA необходимо скачать CUDA Toolkit (на момент данного исследования текущей была версия 5.5), CUDA SDK. SDK можно не скачивать, это всего лишь примеры программ.

Самое главное, необходимо скачать и установить свежие драйвера для видеокарты, поскольку драйвера, поставляемые в комплекте, как правило, устаревшие и не поддерживают текущую версию CUDA. В частности, для CUDA v.5.5 потребовалось скачать драйвера v.335.23.

CUDA Toolkit интегрируется с Visual Stidio, однако есть несколько подводных камней, которые не позволяют сходу начасть работать с данной технологией. В статье [3] приведены полезные советы для решения проблем.

Реализация

Vector - последовательный алгоритм

В качестве массива используется Vector из стандартной библиотеки шаблонов.

using namespace std;

void CalcA_s1(const vector<int> &s,

const vector<int> &sa,

const vector<int> &TetaA,

vector<int> &A,

const int s_dim,

const int a_dim)

{

for (int i = 0; i < a_dim; i++)

{

int sum = 0;

for (int j = 0; j < s_dim; j++)

sum += s[j] * sa[s_dim * i + j];

if (sum > TetaA[i])

A[i] = 1;

else

A[i] = 0;

}

}

Vector - параллельный алгоритм

В качестве контейнера используем по прежнему Vector.Для распараллеливания используем библиотеку PPL (Parallel Patterns Library) входящую в состав Visual C++ 2010.

using namespace Concurrency;

void CalcA_ppl1(const vector<int> &s,

const vector<int> & sa,

const vector<int> &TetaA,

vector<int> &A,

const int s_dim,

const int a_dim

)

{

parallel_for(0, a_dim, [&] (int i){

int sum = 0;

for (int j = 0; j < s_dim; j++)

sum += s[j] * sa[s_dim * i + j];

if (sum > TetaA[i])

A[i] = 1;

else

A[i] = 0;

});

}

Классический массив - последовательный алгоритм

Третью реализацию оформляем в виде класса perceptron в области видимости perceptron_s2, с прицелом переписать на CUDA. Контейнером выступает обычный массив.

namespace perceptron_s2 {

class perceptron {

private:

bool s_foreign;

bool sa_foreign;

bool TetaA_foreign;

bool ra_foreign;

bool TetaR_foreign;

public:

const size_t s_dim;

const size_t a_dim;

const size_t sa_dim;

const size_t r_dim;

const size_t ra_dim;

int *s;

int *sa;

int *TetaA;

int *a;

float *ra;

float *TetaR;

int *r;

perceptron(size_t s_dim, size_t a_dim, size_t r_dim);

perceptron(size_t s_dim, size_t a_dim, size_t r_dim, int *sa, int *TetaA, float *ra, float *TetaR);

perceptron(size_t s_dim, size_t a_dim, size_t r_dim, const vector<int> asa, const vector<int> aTetaA, const vector<float> ara, const vector<float> aTetaR);

void setS(int *as);

void setS(vector<int> &as);

~perceptron();

void show_state();

void calcA();

void calcR();

private:

inline perceptron();

};

}

Приведу только реализации конструктора, принимающего в качестве параметров вектора и функции вычисления CalcA

perceptron::perceptron(

size_t s_dim,

size_t a_dim,

size_t r_dim,

const vector<int> asa,

const vector<int> aTetaA,

const vector<float> ara,

const vector float aTetaR

):

s_dim(s_dim),

a_dim(a_dim),

sa_dim(s_dim * a_dim),

r_dim(r_dim),

ra_dim(r_dim * a_dim)

{

s = new int[s_dim];

s_foreign = false;

sa = new int[sa_dim];

sa_foreign = false;

for(size_t i = 0; i < sa_dim; i++)

sa[i] = asa[i];

TetaA = new int[a_dim];

TetaA_foreign = false;

for(size_t i = 0; i < a_dim; i++)

TetaA[i] = aTetaA[i];

ra = new float[ra_dim];

ra_foreign = false;

for(size_t i = 0; i < ra_dim; i++)

ra[i] = ara[i];

TetaR = new float[r_dim];

TetaR_foreign = false;

for(size_t i = 0; i < r_dim; i++)

TetaR[i] = aTetaR[i];

a = new int[a_dim];

r = new int[r_dim];

}

void perceptron::calcA()

{

for (size_t i = 0; i < a_dim; i++)

{

int sum = 0;

for (size_t j = 0; j < s_dim; j++)

sum += s[j] * sa[s_dim * i + j];

if (sum > TetaA[i])

a[i] = 1;

else

a[i] = 0;

}

}

Последовательная реализация на массивах не отличается от последовательной реализации с использованием stl::Vector. Это даст нам вомзможность сравнить производительность вектора и массива.

CUDA

Реализация CUDA является развитием реализации на массивах. Код оформлен в виде класса perceptron в области видимости perceptron_cuda1, расположен в файле "perceptron_cuda1.cpp".

Код работы с CUDA компилируется отдельным компилятором от NVIDIA, поэтому он вынесен в отдельный файл "perceptron_cuda1_impl.cu". Проект Visual Studio настроен таким образом, что "perceptron_cuda1_impl.cu" компилируется компилятором nvidia во время построения приложения. Все происходит прозрачно и дополнительных хлопот для разработчика не вызывает.

Объявление класса perceptron_cuda1::perceptron незначительно отличается от класса perceptron_s2::perceptron.

Архитектура класса perceptron_cuda1::perceptron следующая. В конструкторе происходит иницализация устройства GPU. Выделяется место в памяти CPU согласно параметрам. В эту память копируются данные из параметров. Выделяется память в GPU и копируются данные в память устройства. Данные - вектора: sa - веса среднего слоя, TetaA - параметры активационной функции среднего слоя, ra - веса выходного слоя и TetaR - параметры активационной функции выходного слоя. Данные массивы сохраняются в памяти устройства GPU и хранятся там все время жизни объекта перцептрона. Во время расчета в процедуру CalcA (расчет среднего слоя) передается массив s. На устройстве GPU выделяется память, копируются данные s, выделяется место под массив среднего слоя A. Запускается на выполнение ядро GPU, которое во множестве потоков производит расчет A. Массив А копируется в память CPU, после чего память GPU с s и A высвобождается.

В классе имеются указатели на sa, TetaA, ra, TetaR в памяти GPU.

Объявление класса:

namespace perceptron_cuda1 {

class perceptron

{

private:

bool s_foreign;

bool sa_foreign;

bool TetaA_foreign;

bool ra_foreign;

bool TetaR_foreign;

int numDevice;

int *dev_sa;

int *dev_TetaA;

int *dev_a;

float *dev_ra;

float *dev_TetaR;

bool isDeviceInitialized;

void releaseDevice();

public:

bool isError;

char *errorMessage;

const size_t s_dim;

const size_t a_dim;

const size_t r_dim;

int *s;

int *sa;

int *TetaA;

float *ra;

float *TetaR;

int *a;

int *r;

perceptron(size_t s_dim, size_t a_dim, size_t r_dim);

perceptron(size_t s_dim, size_t a_dim, size_t r_dim, int *sa, int *TetaA, float *ra, float *TetaR);

perceptron(size_t s_dim, size_t a_dim, size_t r_dim, const vector int asa, const vector int aTetaA, const vector float ara, const vector float aTetaR);

~perceptron();

void show_state();

void initDevice();

void calcA();

void calcA(int *as);

void calcA(vector int &as);

void calcR();

void calc();

void calc(int *as);

void calc(vector int &as);

private: inline perceptron();

};

}

Реализация CalcA:

void perceptron::calcA(int *as)

{

if (!s_foreign)

delete[] s;

s = as;

s_foreign = true;

isError = ! perceptron_cuda1_calcA(s_dim, a_dim, s, dev_sa, dev_TetaA, dev_a, a, &errorMessage);

}

void perceptron::calcA(vector<int> &as)

{

if (s_foreign)

{

s = new int[s_dim];

s_foreign = false;

}

for(size_t i = 0; i < s_dim; i++)

s[i] = as[i];

isError = ! perceptron_cuda1_calcA(s_dim, a_dim, s, dev_sa, dev_TetaA, dev_a, a, &errorMessage);

}

Производится трансляция вызовов процедур в модуле CUDA "perceptron_cuda1_impl.cu". Класс является оберткой CUDA-процедур.

Реализация процедур в модуле "perceptron_cuda1_impl.cu":

#include <stdio.h>

#include iostream

// CUDA runtime

#include <cuda.h>

#include <cuda_runtime.h>

#include <device_launch_parameters.h>

// helper functions and utilities to work with CUDA

#include <helper_cuda.h>

#include <helper_functions.h>

extern "C" bool perceptron_cuda1_initDevice(

size_t s_dim,

size_t a_dim,

size_t r_dim,

int *sa,

int *TetaA,

float *ra,

float *TetaR,

int *numDevice,

int **dev_sa,

int **dev_TetaA,

float **dev_ra,

float **dev_TetaR,

int **dev_a,

char **errorMessage

)

{

*numDevice = 0;

cudaError_t cudaStatus = cudaSuccess;

cudaStatus = cudaSetDevice(*numDevice);

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?";

return false;

}

cudaStatus = cudaMalloc((void**)dev_sa, s_dim * a_dim * sizeof(int));

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMalloc dev_sa failed in initDevice!";

return false;

}

cudaStatus = cudaMalloc((void**)dev_TetaA, a_dim * sizeof(int));

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMalloc dev_TetaA failed in initDevice!";

return false;

}

cudaStatus = cudaMalloc((void**)dev_ra, a_dim * r_dim * sizeof(float));

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMalloc dev_ra failed in initDevice!";

return false;

}

cudaStatus = cudaMalloc((void**)dev_TetaR, r_dim * sizeof(float));

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMalloc dev_TetaR failed in initDevice!";

return false;

}

cudaStatus = cudaMemcpy(*dev_sa, sa, s_dim * a_dim * sizeof(int), cudaMemcpyHostToDevice);

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMemcpy dev_sa failed in calcA!"; return false;

}

cudaStatus = cudaMemcpy(*dev_TetaA, TetaA, a_dim * sizeof(int), cudaMemcpyHostToDevice);

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMemcpy dev_TetaA failed in calcA!";

return false;

}

cudaStatus = cudaMemcpy(*dev_ra, ra, a_dim * r_dim * sizeof(float), cudaMemcpyHostToDevice);

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMemcpy dev_ra failed in calcA!";

return false;

}

cudaStatus = cudaMemcpy(*dev_TetaR, TetaR, r_dim * sizeof(float), cudaMemcpyHostToDevice);

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMemcpy dev_TetaR failed in calcA!";

return false;

}

cudaStatus = cudaMalloc((void**)dev_a, a_dim * sizeof(int));

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMalloc dev_A failed in initDevice!";

return false;

}

return true;

}

__global__ void kernel_calcA(

size_t s_dim,

size_t a_dim,

int *s,

int *sa,

int *TetaA,

int *a

)

{

int tid = threadIdx.x + blockIdx.x * blockDim.x;

while (tid < a_dim)

{

int sum = 0;

for (int j = 0; j < s_dim; j++)

{

sum += s[j] * sa[s_dim * tid + j];

}

if (sum > TetaA[tid])

a[tid] = 1;

else

a[tid] = 0;

tid += blockDim.x * gridDim.x;

}

}

__global__ void kernel_calcR(

size_t a_dim,

size_t r_dim,

int *a,

float *ra,

float *TetaR,

int *r

)

{

int tid = threadIdx.x + blockIdx.x * blockDim.x;

while (tid < r_dim)

{

float sum = 0;

for (int j = 0; j < a_dim; j++)

{

sum += a[j] * ra[a_dim * tid + j];

}

if (sum > TetaR[tid])

r[tid] = 1;

else

r[tid] = -1;

tid += blockDim.x * gridDim.x;

}

}

extern "C" bool perceptron_cuda1_calcA(size_t s_dim, size_t a_dim, const int *s, int *dev_sa, int *dev_TetaA, int *dev_a, int *a, char **errorMessage )

{

const int threadsPerBlock = 256;

const int blocksPerGrid = min(32, (a_dim + threadsPerBlock - 1) / threadsPerBlock);

// printf("blocksPerGrid = %d ",blocksPerGrid);

int *dev_s;

cudaError_t cudaStatus = cudaSuccess;

cudaStatus = cudaMalloc((void**)&dev_s, s_dim * sizeof(int));

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMalloc dev_S failed in calcA!"; return false;

}

cudaStatus = cudaMemcpy(dev_s, s, s_dim * sizeof(int), cudaMemcpyHostToDevice);

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMemcpy s to dev_s failed in calcA!"; goto end;

}

kernel_calcA<<<blocksPerGrid, threadsPerBlock>>>(s_dim, a_dim, dev_s, dev_sa, dev_TetaA, dev_a);

cudaStatus = cudaMemcpy(a, dev_a, a_dim * sizeof(int), cudaMemcpyDeviceToHost);

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMemcpy dev_a to a failed in calcA!";

goto end;

}

end:

cudaFree(dev_s);

return !(cudaStatus != cudaSuccess);

}

extern "C" bool perceptron_cuda1_calcR( size_t a_dim, size_t r_dim, int *dev_a, float *dev_ra, float *dev_TetaR, int *r, char **errorMessage )

{

const int threadsPerBlock = 256;

const int blocksPerGrid = min(32, (r_dim + threadsPerBlock - 1) / threadsPerBlock);

// printf("blocksPerGrid = %d ",blocksPerGrid);

int *dev_r;

cudaError_t cudaStatus = cudaSuccess;

cudaStatus = cudaMalloc((void**)&dev_r, r_dim * sizeof(int));

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMalloc dev_r failed in calcA!";

return false;

}

kernel_calcR<<<blocksPerGrid, threadsPerBlock>>>(a_dim, r_dim, dev_a, dev_ra, dev_TetaR, dev_r);

cudaStatus = cudaMemcpy(r, dev_r, r_dim * sizeof(int), cudaMemcpyDeviceToHost);

if (cudaStatus != cudaSuccess)

{

*errorMessage = "cudaMemcpy dev_r to r failed in calcA!";

goto end;

}

end:

cudaFree(dev_r);

return !(cudaStatus != cudaSuccess);

}

extern "C" void perceptron_cuda1_releaseDevice(int *dev_sa, int *dev_TetaA, int *dev_a, float *dev_ra, float *dev_TetaR)

{

cudaFree(dev_sa);

cudaFree(dev_TetaA);

cudaFree(dev_ra);

cudaFree(dev_TetaR);

cudaFree(dev_a);

}

Анализ производительности

Делаем прогон для следующих размеров перцептронов:

10*15

600*800

1200*1600

4000*6000

Таблица 1. Затраты времени на расчет перцептрона, мсек. ..

где s1 - последовательная реализация, Vector

ppl - параллельная, Vector

s2 - последовательная, массивы

cuda1 - CUDA, массивы

На Рис. 1 приведена диаграмма затрат времени для всех алгоритмов в зависимости от размера перцептрона, на Рис. 2 - тольк для двух последних задач с большим размером перцептрона.