Newton's method for systems of nonlinear equations, scalability3
Основные авторы описания: О.Р.Гирняк, Д.А.Васильков.
Содержание
1 Ссылки
Для исследования масштабируемости алгоритм был реализован нами самостоятельно на библиотеке CUDA 8.0 для языка C++(компилятор nvcc) для Windows.
2 Локальность данных и вычислений
2.1 Локальность реализации алгоритма
2.1.1 Структура обращений в память и качественная оценка локальности
2.1.2 Количественная оценка локальности
3 Масштабируемость алгоритма и его реализации
Масштабируемость алгоритма и его реализаций определяется главным образом масштабируемостью реализации алгоритма решения СЛАУ, используемого для нахождения изменения текущего решения на очередной итерации.
3.1 Масштабируемость алгоритма
3.2 Масштабируемость реализации алгоритма
Использовался домашний компьютер (OS Windows 10 x64, CPU Intel Core i7-2600 3.4 GHz, 8 GB RAM, GPU NVIDIA GTX 550 Ti). GPU поддерживает спецификацию CUDA 2.1 и максимальное число потоков в блоке 1024. Однако физически у нас есть всего 192 CUDA ядра (см. спецификации видеокарты). Все потоки должны быть элементами двумерного квадратного блока. Поэтому исследования масштабируемости были проведены на 4, 16, 36, 64, 100, 144, 196, 256, 324, 400 потоках(на большем количестве потоков проводить исследования бессмысленно ввиду ухудшения производительности). Параметры компиляции:
"nvcc.exe" -gencode=arch=compute_20,code=\"sm_20,compute_20\" --use-local-env --cl-version 2015 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\bin" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include" -G --keep-dir Debug -maxrregcount=0 --machine 32 --compile -cudart static -g -DWIN32 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /FS /Zi /RTC1 /MDd " -o Debug\kernel.cu.obj "kernel.cu"
Решение СЛАУ было реализовано в 2 этапа:
1) Нахождение обратной матрицы с помощью метода Гаусса-Жордана. Для этого было написано ядро для CUDA:
__global__ void gaussjordan(float *A, float *I, int n, int i)
2) Умножение обратной матрицы на вектор правых частей. Для этого было написано ядро для CUDA:
__global__ void multiply_by_vec(float *mat, float *vec, float *out, const int N)
Необходимо отметить, что использование своих ядер для CUDA непременно ведет к снижению производительности. Для улучшения производительности предпочтительно использовать матричные операции из библиотеки cuBLAS. Однако, при таком подходе невозможно явно задать количество потоков - для конкретной видеокарты сама библиотека выбирает, какое количество потоков ей использовать. В связи с этим и были написаны свои ядра, позволяющие задавать количество потоков.
Считалось, что на вход алгоритму были заданы нелинейные функции и их производные. С данной реализацией можно ознакомиться здесь.
Говоря о масштабируемости данного алгоритма, важно отметить следующее:
1. Время вычисления набора функций и набора их производных очень сильно зависит от их сложности и не связано с мощностью алгоритма.
2. Скорость отработки алгоритма от начала до конца может не быть связана с его мощностью: заранее неизвестно начальное приближение, а значит и количество итераций, за которое он сойдется.
Вывод: для оценки масштабируемости необходимо сравнивать одну итерацию алгоритма, и притом, без учета времени вычисления значения функций и их производных. В итоге, приходим к выводу, что целесообразно измерить время выполнения одного решения СЛАУ. Также стоит отметить, что в нашей реализации при хорошем выборе начального приближения алгоритм сходился за 10-20 итераций.
Можно рассмотреть зависимость масштабируемости от одного из двух критериев: размерности задачи или же количества потоков. На графике ниже представлена зависимость времени решения СЛАУ от количества потоков и размерности матрицы.
Как уже было сказано, из-за наличия только 192 физических процессоров, данный алгоритм показывает лучшее время работы при использовании 144 либо 256 потоков. 144 потока (размер блока 12х12) - наиболее близкое количество потоков к 192 физическим процессорам. Хорошее время работы на 256 потоках объясняется тем, что мы используем стандартный размер блока 16х16, а также многие операции деления или умножения на количество потоков или размер блока в этом случае компилятор может заменить на операции побитового сдвига, которые работают быстрее обычных.