|
Архитектура Аудит Военная наука Иностранные языки Медицина Металлургия Метрология Образование Политология Производство Психология Стандартизация Технологии |
|
|
Архитектура Аудит Военная наука Иностранные языки Медицина Металлургия Метрология Образование Политология Производство Психология Стандартизация Технологии |
Сравнение скорости выполнения ⇐ ПредыдущаяСтр 7 из 7
Характеристики компьютера: · Процессор: Intel Core i3-6100 3.7GHz · Видеокарта: GeForce GTX 660 2048Mb Стоит заметить, что скорость работы может различаться среди различных конфигураций компьютера. Основными параметрами, которые определяют скорость работы алгоритма являются процессор и видеокарта, а также пропускная способность материнской платы. Время работы алгоритма фильтрации приведены в таблице 2 (время указано в миллисекундах): Таблица 2 Результаты замеров
В таблице 2 приведены средние результаты, по итогу 100 тестов. Из результатов видно, что на малых значения технология CUDA дает преимущество в 50 раз, а при увеличении размера матрицы, это значение возрастает до 150 раз и больше. Это можно объяснить тем, что для выполнения вычислений на видеокарте требуется их подготовка, копирование в память графического процессора и обратно, поэтому при малых размерах исходных данных, преимущество параллельных вычислений может быть менее заметным. Таким образом, на довольно стандартном размере матрицы 180 х 10 000 в задачах фильтрации, работа с видеокартой позволяет значительно ускорить время работы алгоритма. Также стоит отметить, что исходя из возможностей видеокарты это не предельный размер для работы с технологией CUDA. Так как размер данных не большой, то можно заметить, что потери на подготовку данных для GPU малы в сравнении с общим временем выполнения. В итоге получается матрица, которая представляет собой готовый объект для изучения на предмет выявления дефектов. Большинство шумов исчезло, остались лишь связные наклонные пачки сигналов с наибольшей амплитудой, которые дают объективную картину состояния исследуемого объекта.
Заключение В ходе работы была показана актуальность обработки больших объемов данных с целью выявления дефектов на железнодорожных путях. Изучена современная технология CUDA для решения поставленной задачи фильтрации. Предложен алгоритм и методы его оптимизации. На основе полученных тестов было продемонстрировано, что работа с видеокартой может значительно ускорить выполнение алгоритма за счет большого количества параллельных потоков. В целом, выполнение задач с использованием видеокарты в современном информационном обществе находит все больше применения. Работа с большими данными требует больших вычислительных мощностей, которые может обеспечить видеокарта, либо несколько процессоров. На основе полученных тестов можно убедиться, что видеокарта дает заметное преимущество в обработке значительного числа информации. Также стоит заметить, что технология CUDA хорошо справляется с большим объемом арифметических операции, но менее применима к структурам, где требуется множество ветвлений.
Список литератур ы 1. Боресков А.В. Параллельные вычисления на GPU. Архитектура и программная модель CUDA: Учебное пособие. Издательство Московского университета, 2012. 86с. 2. Боресков А.В., Харламов А.А. Основы работы с технологией CUDA. 2010. 40 - 56с. 3. Сандерс Д., Кэндрот Эю. NVIDIA CUDA в примерах: введение в программирование графических процессоров – ДМК Пресс -2011. 4. Седжвик Р. Фундаментальные алгоритмы на C++. М., СПб., Киев: DiaSoft, 2001. 152с. 5. Перепёлкин Е.Е., Садовников Б.И., Иноземцева Н.Г. Вычисления на графических процессорах (GPU) в задачах математической и теоретической физики. Издательство: URSS, 2014. 15с. 6. Марков А.А., Кузнецова Е.А. Дефектоскопия рельсов. Формирование и анализ сигналов. Санкт-Петербург, Ультра Принт, 2014. 212с. 7. Тарабрин В.Ф., Зверев А. В., Горбунов О.Е., Кузьмин Е.В. О фильтрации данных..., В мире неразрушающего контроля 2 (64) июнь 2014, 5 - 9 с. 8. Официальная документация по технологии CUDA: http: //www.nvidia.ru. 1веб-страница. URL: http: //www.nvidia.ru/object/cuda_home_new_ru.html. 9. Описание технологии CUDA: http: //bourabai.kz. 1веб-страница. URL: http: //bourabai.kz/graphics/cuda.htm. 10. Неграфические вычисления CUDA: https: //www.ixbt.com. 1веб-страница. URL: https: //www.ixbt.com/video3/cuda-1.shtml. 11. Параллельные вычисления: https: //studwood.ru. 1веб-страница. URL: https: //studwood.ru/1622722/informatika/tehnologiya_nvidia_cuda. 12. Cuda - Основные примеры: http: //www.ksmlab.ru. 1веб-страница. URL: http: //www.ksmlab.ru/page/cuda-osnovy-primery. 13. Дорофеев А.Л., Кузаманов Ю.Г. Электромагнитная дефектоскопия. — 2-е изд., доп. и перераб. -М.: Машиностроение, 1980. —232 е.: ил. 14. Горделий В.И. Состояние и перспективы развития средств неразрушающего контроля рельсов // Нефть, газ, промышленность №6, 2005. 15. Правила технической эксплуатации железных дорог Российской Федерации. - министерство путей сообщения Российской Федерации. — М.: «Техинформ», - 2000г. 114 - 190 с. Приложение А. Листинг программы по технологии CUDA В данном файле представлен код, который выполняется на видеокарте, его можно отличить по приставке __device__ или __global__. Остальные же методы служат для копирования данных на видеокарту и обратно, выделение памяти, разметки процессов и непосредственно запуска ядра видеокарты.
Файл kernel. cu // Библиотеки для работы с CUDA Runtime API 8.0 #include " cuda_runtime.h" #include " device_launch_parameters.h"
#include < stdio.h> #include < iostream>
/* * Проверка на принадлежность элемента к массиву, метод выполняется на видеокарте */ __device__ bool contains(unsigned short int* indexes, int n, unsigned short index) { for (int i = 0; i < n; ++i) { if (indexes[i] == index) { return true; } } return false; }
/* * Фильтрация данных по столбцам, метод запускается на видеокарте */ __global__ void verticalFilteringWithSearching(unsigned short int* dev, int const n, int const height, int const sizeOfMask, int const keepCount) { int const startIndex = threadIdx.x + threadIdx.y * sizeOfMask + blockDim.x * blockIdx.x;
unsigned short int* maxIndexes = new unsigned short[keepCount];
for (int i = 0; i < keepCount; i++) maxIndexes[i] = -1;
for (int k = 0; k < keepCount; ++k) { short int maxValue = -1; int maxIndex = -1; for (int i = startIndex; i < startIndex + sizeOfMask; i += n) { unsigned short value = dev[i]; if (value > = maxValue & & ! contains(maxIndexes, keepCount, i)) { maxValue = value; maxIndex = i; } } maxIndexes[k] = maxIndex; }
for (int i = startIndex; i < startIndex + sizeOfMask; i += n) { if (! contains(maxIndexes, keepCount, i)) dev[i] = 0; } }
/* * Фильтрация данных по строкам, метод запускается на видеокарте */ __global__ void horizontalFilteringWithSearching(unsigned short int* dev, int const n, int const sizeOfMask, int const keepCount) { int const startIndex = n * blockIdx.y + sizeOfMask * threadIdx.x;
unsigned short int* maxIndexes = new unsigned short[keepCount];
for (int i = 0; i < keepCount; i++) maxIndexes[i] = -1;
for (int k = 0; k < keepCount; ++k) { short int maxValue = -1; int maxIndex = -1; for (int i = startIndex; i < sizeOfMask + startIndex; ++i) { unsigned short value = dev[i]; if (value > = maxValue & & ! contains(maxIndexes, keepCount, i)) { maxValue = value; maxIndex = i; } } maxIndexes[k] = maxIndex; }
for (int i = startIndex; i < startIndex + sizeOfMask; ++i) { if (! contains(maxIndexes, keepCount, i)) dev[i] = 0; } } /* * Метод фильтрации по строкам с использование технологии CUDA */ void cudaHorizontalFilter(short int* data, int const n, int const height, int const sizeOfMask, int const keepCount) { cudaError_t cudaStatus; unsigned short int* dev = nullptr; int numBytes = n * height * sizeof(short int); if (n % sizeOfMask) { std:: cout < < " WARNING: size of mask is not the multiplicity for size line"; } std:: cout < < " Start cuda horizontal filter\n";
cudaStatus = cudaSetDevice(0); if (cudaStatus! = cudaSuccess) { fprintf(stderr, " cudaSetDevice failed! Do you have a CUDA-capable GPU installed? " ); } cudaStatus = cudaMalloc((void**)& dev, numBytes); if (cudaStatus! = cudaSuccess) { fprintf(stderr, " cudaMalloc failed! " ); }
dim3 threads = dim3(n / sizeOfMask, 1); dim3 blocks = dim3(1, height);
cudaMemcpy(dev, data, numBytes, cudaMemcpyHostToDevice);
horizontalFilteringWithSearching< < < blocks, threads > > > (dev, n, sizeOfMask, keepCount);
cudaMemcpy(data, dev, numBytes, cudaMemcpyDeviceToHost);
cudaFree(dev); }
/* * Метод фильтрации по столбцам с использование технологии CUDA */ void cudaVerticalFilter(short int* data, int const n, int const height, int const sizeOfMask, int const keepCount) { unsigned short int* dev = nullptr; int numBytes = n * height * sizeof(short int); if (height % sizeOfMask) { std:: cout < < " WARNING: size of mask is not the multiplicity for size line"; } std:: cout < < " Start cuda vertical filter\n";
cudaMalloc((void**)& dev, numBytes);
dim3 threads = dim3(n / 100, height / sizeOfMask); dim3 blocks = dim3(100, 1);
cudaMemcpy(dev, data, numBytes, cudaMemcpyHostToDevice); verticalFilteringWithSearching < < < blocks, threads > > > (dev, n, height, sizeOfMask*n, keepCount);
cudaMemcpy(data, dev, numBytes, cudaMemcpyDeviceToHost);
cudaFree(dev); }
/* * Фильтрация данных по диагонали, метод запускается на видеокарте */ __global__ void diagonalFilteringWithSearching(unsigned short int* dev, int const n, int const sizeOfMask, int const keepCount) { int const startIndex = sizeOfMask * threadIdx.x;
unsigned short int* maxIndexes = new unsigned short[keepCount];
for (int i = 0; i < keepCount; i++) maxIndexes[i] = -1;
int maxIndex = -1; for (int k = 0; k < keepCount; ++k) { short int maxValue = -1; for (int i = startIndex; i < sizeOfMask + startIndex; ++i) { unsigned short value = dev[i]; if (value > = maxValue & & ! contains(maxIndexes, keepCount, i)) { maxValue = value; maxIndex = i; } } if (startIndex + sizeOfMask + 1 < n) { int maxIndexNeighbor = -1; short int maxValueNeighbor = -1; for (int i = startIndex; i < sizeOfMask + startIndex; ++i) { unsigned short value = dev[i]; if (value > = maxValueNeighbor & & ! contains(maxIndexes, keepCount, i)) { maxValueNeighbor = value; maxIndexNeighbor = i; } } if (abs(maxIndexNeighbor - maxIndex) < 10) { maxIndexes[k] = maxIndex; } } else { maxIndexes[k] = maxIndex; } }
for (int i = startIndex; i < startIndex + sizeOfMask; ++i) { if (! contains(maxIndexes, keepCount, i)) dev[i] = 0; } }
/* * Метод фильтрации по диагонали с использование технологии CUDA */ void cudaDiagonalFilter(short int* data, int const n, int const height, int const sizeOfMask, int const keepCount) { cudaError_t cudaStatus; unsigned short int* dev = nullptr; int numBytes = n * height * sizeof(short int); if (n % sizeOfMask) { std:: cout < < " WARNING: size of mask is not the multiplicity for size line"; } std:: cout < < " Start cuda diagonal filter\n";
cudaStatus = cudaSetDevice(0); if (cudaStatus! = cudaSuccess) { fprintf(stderr, " cudaSetDevice failed! Do you have a CUDA-capable GPU installed? " ); }
cudaStatus = cudaMalloc((void**)& dev, numBytes); if (cudaStatus! = cudaSuccess) { fprintf(stderr, " cudaMalloc failed! " ); }
dim3 threads = dim3(n / sizeOfMask, 1); dim3 blocks = dim3(1, 1);
cudaMemcpy(dev, data, numBytes, cudaMemcpyHostToDevice);
diagonalFilteringWithSearching < < < blocks, threads > > > (dev, n, sizeOfMask, keepCount);
cudaMemcpy(data, dev, numBytes, cudaMemcpyDeviceToHost);
cudaFree(dev); }
/* * Получение информации о всех видеокартах компьютера */ void getInfo() { int deviceCount; cudaDeviceProp devProp;
cudaGetDeviceCount(& deviceCount);
printf(" Found %d devices\n", deviceCount);
for (int device = 0; device < deviceCount; device++) { cudaGetDeviceProperties(& devProp, device);
printf(" Device %d\n", device); printf(" Compute capability: %d.%d\n", devProp.major, devProp.minor); printf(" Name : %s\n", devProp.name); printf(" Total Global Memory: %d\n", devProp.totalGlobalMem); printf(" Shared memory per block: %d\n", devProp.sharedMemPerBlock); printf(" Registers per block: %d\n", devProp.regsPerBlock); printf(" Warp size : %d\n", devProp.warpSize); printf(" Max threads per block: %d\n", devProp.maxThreadsPerBlock); printf(" Total constant memory: %d\n", devProp.totalConstMem);
} }
Приложение Б. Листинг программы без технологии CUDA В данном файле представлен код, который считывает исходную матрицу значений, затем запускает фильтрацию на процессоре и видеокарте с замерами времени соответственно. Размер исходные данные задается заранее и не изменяется в процессе фильтрации. Методы для работы с технологией CUDA импортируются как внешние зависимости, поэтому компилируются отдельно от основного файла. После выполнения фильтрации данные сохраняются в файл, для каждого способа отдельно. Файл Main.cpp #include < iostream> #include < fstream> #include < sstream> #include < iterator> #include < vector> #include < ctime> #include < iomanip> #include < thread> #include < chrono>
// Подключение функций работающих с CUDA extern void cudaHorizontalFilter(short int* data, int const n, int const height, int const sizeOfMask, int const keepCount); extern void cudaVerticalFilter(short int* data, int const n, int const height, int const sizeOfMask, int const keepCount); extern void diagonalFilteringWithSearching(unsigned short int* dev, int const n, int const sizeOfMask, int const keepCount); extern void getInfo();
// Имя файла для чтения данных std:: string const pathToData = " smp1.txt"; int const countLines = 180; int const countColumns = 10000;
// Переменные для выбора способа фильтрации std:: string onlyCPU = " cpu"; std:: string onlyGPU = " gpu"; std:: string any = " any";
int getRightSide(int sizeOfMask, int i, int j, std:: vector< std:: vector< short int> > & data, int& keepCount);
int getDownSide(int sizeOfMask, int i, int j, std:: vector< std:: vector< short int> > & data, int& keepCount);
std:: vector< short int> getKeepIndexesHorizontal(int keepCount, int i, int j, int rightSide, std:: vector< std:: vector< short int> > & data);
std:: vector< short int> getKeepIndexesVertical(int keepCount, int i, int j, int downSide, std:: vector< std:: vector< short int> > & data); std:: vector< short int> getKeepIndexesDiagonal(int i, int j, int rightSide, std:: vector< std:: vector< short int> > & data); /* * Получение данных из файла */ std:: vector< std:: vector< short int> > getDatafromFile() { std:: vector< std:: vector< short int> > data; std:: ifstream reader(pathToData, std:: ios_base:: in); std:: string str; if (! reader.is_open()) { std:: cerr < < " Error: file " < < pathToData < < " is not open\n"; return data; } int counter = 0; while (getline(reader, str) & & counter! = countLines) { counter++; std:: istringstream iss(str); std:: vector< std:: string> results((std:: istream_iterator< std:: string> (iss)), std:: istream_iterator< std:: string> ()); std:: vector< short int> dataLine; if (results.size() < countColumns) std:: cerr < < " Error: size of columns is less then need\n";
for (int i = 0; i < countColumns; ++i) { dataLine.push_back((short int)std:: stoi(results[i])); } data.push_back(dataLine); } reader.close(); return data; }
/* * Сохранение данных после фильтрации на процессоре */ void saveDataToFile(std:: string filename, std:: vector< std:: vector< short int> > data) { std:: ofstream file(filename, std:: ios:: out); if (! file.is_open()) { std:: cerr < < " Error: file " < < filename < < " is not open\n"; return; }
for (int i = 0; i < data.size(); ++i) { for (int j = 0; j < data[0].size(); ++j) { file < < std:: setw(4) < < data[i][j] < < " "; } file < < std:: endl; }
file.close(); }
/* * Сохранение данных после фильтрации на видеокарте */ void saveDataToFile(std:: string filename, short int* data, int n, int height) { std:: ofstream file(filename, std:: ios:: out); if (! file.is_open()) { std:: cerr < < " Error: file " < < filename < < " is not open\n"; return; }
for (int i = 0; i < height; ++i) { for (int j = 0; j < n; ++j) { file < < std:: setw(4) < < data[i * n + j] < < " "; } file < < " \n"; }
file.close(); }
/* * Проверка на принадлежность элемента к массиву */ bool contains(std:: vector< short int> a, int source) { for (int i = 0; i < a.size(); ++i) { if (a[i] == source) { return true; } } return false; }
/* * Фильтрация данных по столбцам на процессоре */ void verticalFiltering(std:: vector< std:: vector< short int> > & data, int sizeOfMask, int keepCount) { for (int i = 0; i < data[0].size(); ++i) { for (int j = 0; j < data.size(); j += sizeOfMask) { int downSide = getDownSide(sizeOfMask, i, j, data, keepCount); std:: vector< short int> keepsIndexes = getKeepIndexesVertical(keepCount, i, j, downSide, data);
for (int m = j; m < downSide; ++m) { if (! contains(keepsIndexes, m)) { data[m][i] = 0; } } } } }
/* * Фильтрация данных по строкам на процессоре */ void horizontalFiltering(std:: vector< std:: vector< short int> > & data, int sizeOfMask, int keepCount) { for (int i = 0; i < data.size(); ++i) { for (int j = 0; j < data[i].size(); j += sizeOfMask) { int rightSide = getRightSide(sizeOfMask, i, j, data, keepCount); std:: vector< short int> keepsIndexes = getKeepIndexesHorizontal(keepCount, i, j, rightSide, data);
for (int m = j; m < rightSide; ++m) { if (! contains(keepsIndexes, m)) { data[i][m] = 0; } } } } }
/* * Фильтрация данных по диагонали на процессоре */ void diagonalFiltering(std:: vector< std:: vector< short int> > & data, int sizeOfMask, int minSignal) { for (int i = 0; i < data.size(); ++i) { for (int j = 0; j < data[i].size(); j += sizeOfMask) { int notUsed = 0; int rightSide = getRightSide(sizeOfMask, i, j, data, notUsed); std:: vector< short int> keepsIndexes = getKeepIndexesDiagonal(i, j, rightSide, data);
for (int m = j; m < rightSide; ++m) { if (! contains(keepsIndexes, m)) { data[i][m] = 0; } } } } }
/* * Получение списка индексов сигналов в строке, которые нужно сохранить. Остальные обнуляются. */ std:: vector< short int> getKeepIndexesHorizontal(int keepCount, int i, int j, int rightSide, std:: vector< std:: vector< short int> > & data) { std:: vector< short int> keepsIndexes; for (int l = 0; l < keepCount; ++l) { int max = -1; int indexOfMax = 0; for (int k = j; k < rightSide; ++k) { if (data[i][k] > = max & & ! contains(keepsIndexes, k)) { max = data[i][k]; indexOfMax = k; } } keepsIndexes.push_back(indexOfMax); } return keepsIndexes; }
/* * Получение списка индексов сигналов в столбце, которые нужно сохранить. Остальные обнуляются. */ std:: vector< short int> getKeepIndexesVertical(int keepCount, int i, int j, int downSide, std:: vector< std:: vector< short int> > & data) { std:: vector< short int> keepsIndexes; for (int l = 0; l < keepCount; ++l) { int max = -1; int indexOfMax = 0; for (int k = j; k < downSide; ++k) { if (data[k][i] > = max & & ! contains(keepsIndexes, k)) { max = data[k][i]; indexOfMax = k; } } keepsIndexes.push_back(indexOfMax); } return keepsIndexes; }
/* * Получение списка индексов сигналов в строке, которые нужно сохранить. Остальные обнуляются. */ std:: vector< short int> getKeepIndexesDiagonal(int i, int j, int rightSide, std:: vector< std:: vector< short int> > & data) { std:: vector< short int> keepsIndexes; int indexOfMax = 0; for (int l = 0; l < 1; ++l) { int max = -1; for (int k = j; k < rightSide; ++k) { if (data[i][k] > = max & & ! contains(keepsIndexes, k)) { max = data[i][k]; indexOfMax = k; } } } if (data[i].size() > j + 1) { int indexOfMaxNeighbor = 0; for (int k = j; k < rightSide; ++k) { int max = -1; if (data[i][k] > = max & & ! contains(keepsIndexes, k)) { max = data[i][k]; indexOfMaxNeighbor = k; } } if (abs(indexOfMaxNeighbor - indexOfMax) < 10) { keepsIndexes.push_back(indexOfMax); } } else { keepsIndexes.push_back(indexOfMax); } return keepsIndexes; }
/* * получение правой границы маски в фильтрации на процессоре */ int getRightSide(int sizeOfMask, int i, int j, std:: vector< std:: vector< short int> > & data, int& keepCount) { int rightSide; if (j + sizeOfMask < data[i].size()) { rightSide = j + sizeOfMask; if (rightSide + sizeOfMask > data[i].size()) { keepCount *= (int)(sizeOfMask + data[i].size() - rightSide) / sizeOfMask; rightSide = (int)data[i].size(); } } else { rightSide = (int)data[i].size(); } return rightSide; }
/* * Получение нижней границы маски в филтрации на процессоре */ int getDownSide(int sizeOfMask, int i, int j, std:: vector< std:: vector< short int> > & data, int& keepCount) { int downSide; if (j + sizeOfMask < data.size()) { downSide = j + sizeOfMask; if (downSide + sizeOfMask > data.size()) { keepCount *= (int)(sizeOfMask + data.size() - downSide) / sizeOfMask; downSide = (int)data.size(); } } else { downSide = (int)data.size(); } return downSide; }
/* * Фильтрация с использование технологии CUDA */ void cudaHandling(short int* data) { cudaHorizontalFilter(data, countColumns, countLines, 20, 5); cudaHorizontalFilter(data, countColumns, countLines, 200, 20); cudaVerticalFilter(data, countColumns, countLines, 20, 5); cudaDiagonalFilter(data, countColumns, countLines, 20, 5); }
/* * Фильтрация данных с помощью процессора */ void cpuHandling(std:: vector< std:: vector< short int> > & data) { horizontalFiltering(data, 20, 5); horizontalFiltering(data, 200, 20); verticalFiltering(data, 20, 5); diagonalFiltering(data, 20, 5); }
/* * Старт программы, чтение данных из файла, запуск методов фильтрации с сохранение данных в файлы */ int main() { time_t start, stop; double timeCPU = 0, timeGPU = 0; std:: string mode = any; std:: cout < < " ------------START READING DATA------------------------------\n";
std:: vector< std:: vector< short int> > vec = getDatafromFile(); if (vec.size()! = countLines || vec[0].size()! = countColumns) { std:: cout < < " ------------ERROR SIZE IS WRONG------------------------------\n"; return 1; } short int* data = new short int[countColumns * countLines]; for (int i = 0; i < countLines; ++i) { for (int j = 0; j < countColumns; ++j) { data[i * countColumns + j] = vec[i][j]; } }
std:: cout < < " ------------DATA SUCCESSFULLY OBTAINED------------------------------\n"; if (mode == onlyCPU || mode == any) { std:: cout < < " ------------CPU START------------------------------\n"; time(& start); timeCPU = clock(); cpuHandling(vec); time(& stop); timeCPU = clock() - timeCPU; saveDataToFile(" cpuresult.txt", vec); std:: cout < < " ------------CPU FINISH------------------------------\n"; }
if (mode == onlyGPU || mode == any) { std:: cout < < " ------------GPU START------------------------------\n"; time(& start); timeGPU = clock(); cudaHandling(data); time(& stop); timeGPU = clock() - timeGPU; saveDataToFile(" gpuresult.txt", data, countColumns, countLines); std:: cout < < " ------------GPU FINISH------------------------------\n"; }
std:: cout < < " CPU time: " < < timeCPU < < " ms\n" < < " GPU time: " < < timeGPU < < " ms\n";
std:: cout < < " ------------FINISH PROGRAM----------------------------------\n"; free(data); return 0; }
|
Последнее изменение этой страницы: 2019-03-20; Просмотров: 332; Нарушение авторского права страницы