Алгоритм параллельной агрегации данных для визуализации данных о вербальном…
9
Так, при
W
= 1920 количество элементов, которые могут быть об-
работаны
GPU
,
9
1920·128·4096 10
N
. Пусть
ShM
— размер раз-
деляемой памяти, доступной одному блоку потоков. Тогда, макси-
мальное количество элементов
Npb
, обрабатываемых одним блоком,
может быть вычислено как
/ (2( / 8)) 4 / .
Npb ShM B
ShM B
(10)
При
16Кб,
ShM
4 ·16384 / 8 8129
Npb
элементов. Таким об-
разом, в цикле необходимо подгружать элементы из глобальной па-
мяти (
global memory
)
128·4096 / 8129 65
раз. Следует заметить, что
доступ к глобальной памяти достигает 80 Гб/с, что позволяет быстро
обращаться к входным данным из процессов
GPU
[12].
Ниже приведен листинг функции вычисления экстремумов для
блока данных
i
s
. При правильном подборе размера блоков
(
blockSize
)
и размера сетки блоков (
blockDim
) инструкции для всех
блоков выполняются одновременно.
// шаблонная функция вычисления минимума и максимума
блока данных
// в качестве типа T может быть int, byte, float и другие
template <class T, unsigned int blockSize>
__global__ void minmaxReduction(T * input, T * out_min,
T * out_max, int len) {
T* sMin = SharedMemory<T>();
// объявляем extern
__shared__ переменные
T* sMax = &sMin[blockSize];
…
// инициализация параметров
while (i < len)
{
// загружаем минимумы/максимумы в переменную sMin
(sMax)
// проверяем граничные условия
if ( i + blockDim.x < len)
sMin[t] = min(sMin[t], min(input[i], input[i + blockDim.x]));
else
sMin[t] = min(sMin[t], input[i]);
i += gridSize; // подгружаем следующий блок данных
}
__syncthreads();
// синхронизируем потоки внутри блока
