;
}
}
//Синхронизация нужна для того чтобы убедиться
//что все нити закончили работу с частью вектора
//потому что следующая операция - перезапись разделяемой памяти
__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[j + k * THREAD_PER_BLOCK] * shared_vector;
}
}
//Синхронизация нужна для того чтобы убедиться
//что все нити закончили работу с частью вектора
//потому что следующая операция - перезапись разделяемой памяти
__syncthread();
}
if ( i < N ) {
С = shared_c[threadId.x];
}
__syncthread();
Модификация кода хоста в этой лабораторной работе не требуется
Для анализа эффективности доступа к разделяемой памяти используйте утилиту nvprof. События для профилирования: shared_ld_bank_conflict – количество конфликтов банков памяти при считывании данных, shared_st_bank_conflict– количество конфликтов банков памяти при записи данных. И метрики: shared_efficiency – эффективность использования пропускной способности шины данных разделяемой памяти, shared_load_transactions_per_request – количество транзакций при каждом запросе к разделяемой памяти.
По аналогии оптимизируйте алгоритм умножения вектора на матрицу.
СКАЧАТЬ
