Программирование графических процессоров Лабораторная 2

Задание
1. Прочитайте главу из теоретического материала "Разделяемая память" и ответьте на контрольные вопросы (ответы на контрольные вопросы не нужно включать в отчёт по лабораторной работе).
2. Оптимизируйте алгоритмы, реализованные в лабораторной работе №1 при помощи разделяемой памяти.
3. Постройте графики зависимости времени выполнения алгоритма от размера матрицы и вектора (Размеры матрицы 1000x500, 1000x1000, 1500x1000, 2000x1000, 2000x1500, 2500x1500, 2500x2000).
4. Проанализируйте, реализованные алгоритмы при помощи утилиты nvprof на эффективность доступа к глобальной памяти.
Методические указания по выполнению лабораторной работы
Разделяемая память – это своего рода кэш. Ускорения от использования разделяемой памяти можно достичь только если к каким-то данным происходит многократное обращение. Тогда переместив их в разделяемую память из глобальной можно сократить время затраченное на обращение в память за счёт высокой скорости разделяемой памяти.
Рассмотрим схему параллельного умножения матрицы на вектор.

Рис. 1 – схема умножения матрицы на вектор
Каждая нить берёт по одной строке матрицы и умножает попарно элементы строки на элементы вектора. Все нити используют один и тот же вектор. Скопировав вектор в разделяемую память можно получить ускорения выполнения алгоритма. Но проблема разделяемой памяти в том, что её объём очень мал по сравнению с объёмом глобальной памяти. Поэтому нужно предусмотреть ситуацию когда вектор полностью в память не помещается.
Учитывая, что разделяемой памяти может не хватить нужно модифицировать алгоритм таким образом чтобы часть вектора копировалась в разделяемую память. Нити умножали часть строки на часть вектора, расположенного в разделяемой памяти. Затем копировали следующую часть вектора в разделяемую память и умножали элементы следующей части строки и т.д. как показано на рисунке 2.

Рис. 2 – умножение нитью строки на вектор с копированием частей вектора в разделяемую память
Копирование частей вектора можно производить параллельно, каждая нить может копировать в разделяемую память один или несколько элементов. Таким образом псевдокод алгоритма будет выглядеть следующим образом:
//Выделяем разделяемую память
//Количество элементов будет равно количеству нитей в блоке.
__shared__ shared_vector[THREAD_PER_BLOCK];

//k - номер части вектора. +1 потому что M не обязательно делится нацело
for(k = 0; k < (M / THREAD_PER_BLOCK) + 1; ++k) {
//p - реальный номер элемента вектора
//который будет скопирован нитью в разделяемую память
p = threadIdx.x + THREAD_PER_BLOCK * k;
//Проверка на выход за пределы вектора
if(p < M) {
//Копирование из глобальной памяти в разделяемую
//Т.к. размер разделяемой памяти равен размеру блока
//Каждая нить может скопировать по одному элементу
 shared_vector[threadIdx.x] = V[threadIdx.x + THREAD_PER_BLOCK * k];
}
//Синхронизация. Нужно дождаться пока все нити скопируют свои элементы
//Потому что каждая нить использует скопированную часть вектора полностью
__syncthread();
//i - глобальный номер нити, N - количество строк в матрице
//Проверка не выходим ли за границы матрицы
if (i < N) {
//Умножение части вектора на часть строки
//j - порядковый номер элемента в части вектора или строки матрицы
//чтобы получить реальный номер элемента матрицы
//нужно прибавить j количеству полностью обработанных элементов
 for(j = 0; j < THREAD_PER_BLOCK; ++j) {
 C[i] += A[i][j + k * THREAD_PER_BLOCK] * shared_vector[j];
}
}
//Синхронизация нужна для того чтобы убедиться
//что все нити закончили работу с частью вектора
//потому что следующая операция - перезапись разделяемой памяти
__syncthread();
}
В данном коде можно сделать ещё одну оптимизацию – перенести операции над вектором C в разделяемую память. И скопировать результат в глобальную только после того как результат будет посчитан. С учётом этого алгоритм будет выглядеть следующим образом:
//Выделяем разделяемую память
//Количество элементов будет равно количеству нитей в блоке.
__shared__ shared_vector[THREAD_PER_BLOCK];
__shared__ shared_c[THREAD_PER_BLOCK];

shared_c[threadId.x] = 0;
__syncthread();

//k - номер части вектора. +1 потому что M не обязательно делится нацело
for(k = 0; k < (M / THREAD_PER_BLOCK) + 1; ++k) {
//p - реальный номер элемента вектора
//который будет скопирован нитью в разделяемую память
p = threadIdx.x + THREAD_PER_BLOCK * k;
//Проверка на выход за пределы вектора
if(p < M) {
//Копирование из глобальной памяти в разделяемую
//Т.к. размер разделяемой памяти равен размеру блока
//Каждая нить может скопировать по одному элементу
 shared_vector[threadIdx.x] = V[threadIdx.x + THREAD_PER_BLOCK * k];
}
//Синхронизация. Нужно дождаться пока все нити скопируют свои элементы
//Потому что каждая нить использует скопированную часть вектора полностью
__syncthread();
//i - глобальный номер нити, N - количество строк в матрице
//Проверка не выходим ли за границы матрицы
if (i < N) {
//Умножение части вектора на часть строки
//j - порядковый номер элемента в части вектора или строки матрицы
//чтобы получить реальный номер элемента матрицы
//нужно прибавить j количеству полностью обработанных элементов
 for(j = 0; j < THREAD_PER_BLOCK; ++j) {
 shared_c[threadIdx.x] += A[i][j + k * THREAD_PER_BLOCK] * shared_vector[j];
}
}
//Синхронизация нужна для того чтобы убедиться
//что все нити закончили работу с частью вектора
//потому что следующая операция - перезапись разделяемой памяти
__syncthread();
}
if ( i < N ) {
 С[i] = shared_c[threadId.x];
}
__syncthread();
Модификация кода хоста в этой лабораторной работе не требуется
Для анализа эффективности доступа к разделяемой памяти используйте утилиту nvprof. События для профилирования: shared_ld_bank_conflict – количество конфликтов банков памяти при считывании данных, shared_st_bank_conflict– количество конфликтов банков памяти при записи данных. И метрики: shared_efficiency – эффективность использования пропускной способности шины данных разделяемой памяти, shared_load_transactions_per_request – количество транзакций при каждом запросе к разделяемой памяти.
По аналогии оптимизируйте алгоритм

зачет