Вторая контрольная работа (1127792), страница 3

Просмтор этого файла доступен только зарегистрированным пользователям. Но у нас супер быстрая регистрация: достаточно только электронной почты!

Текст из файла (страница 3)

Одним из методов уменьшения потерь на выполнение команд переходов, как условных, так и безусловных, является использования кэш-памяти прогнозирования переходов, называемой таблицей, буфером целевых адресов переходов - BTB (Branch Target Buffer). Буфер содержат информацию о командах переходов, последних по времени выполнения программы. Если выполняемая команда перехода не представлена в таблице (пусть её объем равен 2n ), то в i-строку таблицы записывается TAG и <целевые адреса команды>. Номер строки таблицы – i есть значения n-младших разрядов адреса команды перехода (счетчика команд), а старшие разряды адреса этой команды перехода заносятся в поле TAG. Поле <целевые адреса команды> может состоять из исполнительного адреса перехода, который формируется при выполнении команды перехода. Если при выборки очередной команды выясняется факт нахождения ее в таблице ВТВ, тогда для команды безусловного перехода сразу известен исполнительный адрес перехода еще до дешифрации выбранной команды. Выбранная команда игнорируется, в счетчик команд заносится исполнительный адрес перехода из таблицы. Более того, в поле <целевые адреса команды> можно хранить код команды, на которую указывает исполнительный адрес перехода. Для обработки команд условного перехода ВТВ используется совместно с аппаратом динамического предсказания переходов.

Динамическое предсказание переходов опирается на предысторию выполнения вычислительного процесса - для каждого конкретного случая перехода накапливается статистика поведения.

Аппаратный механизм учета вероятности перехода (динамический) состоит из блока предсказания переходов. Этот блок, кроме учета статически определенных предпочтений для ветвлений, использует таблицу переходов, в которой хранится история переходов для каждого (в рамках объема таблицы) перехода программы. Большинство современных микропроцессоров обещают точность предсказаний переходов этим способом выше 90%.

Для динамического прогнозирования переходов аппаратура собирает статистику переходов, которая помещается в – BHT (Branch History Table). Реализация этой схемы имеет ряд красивых архитектурных решений. Таблица ВНТ содержит статистику о результатах вычисления условных выражений: каждая её строка соответствует выполненным условным операторам и она содержит информацию о том, как выполнялись переходы при предыдущих обращениях к данному оператору. Такой информацией в общем случае является значение n-битового счетчика в строке таблицы и принципиальная схема прогнозирования перехода для текущего условного оператора такова.

1. Если значение счетчика больше нуля, то переход прогнозируется как выполняемый. Если направление перехода предсказано верно, к значению счетчика прибавляется единица, иначе значение счетчика уменьшается на единицу.

2. Если значение счетчика меньше нуля, то переход прогнозируется как невыполняемый. Если направление перехода предсказано верно, из значения счетчика вычитается единица, при неверном прогнозе - к значению счетчика единица добавляется.

Для упрощения работы со счетчиком в качестве “нуля”, нейтральной точки отсчета, обычно выбирается точка из середины интервала значений счетчика (например, значение 2n – 1). При этом отказ от работы с отрицательными числами не отменяет процедуры проверки выхода значения счетчика за пределы интервала.

Исследования n-битовых схем прогнозирования на тестовых задачах SPEC показало, что наилучшие результаты прогнозирования обеспечивают трех-битовые счетчики, причем результаты двух-битовых отличаются от трех-битовых незначительно. Большинство схем прогнозирования используют двух-битовые счетчики, причем для работы с ними используется алгоритм Смита. В соответствии с этим алгоритмом, значения счетчиков определяют четыре вероятности: ветвь часто исполняется (strongly taken, код – 11), ветвь выполняется (taken, код – 10), ветвь не выполняется (not taken, код -01), ветвь часто не выполняется (strongly not taken, код – 00). В этой схеме прогнозирования переход прогнозируется, если первый, левый бит равен единице. Правый бит в этой схеме показывает результат вычисления условия на предыдущем шаге.

Другой проблемой является выбор размера таблицы ВНТ. Статистические исследования показывают, что таблица, состоящая из 4096 строк, дает высокую точность прогнозов и на ряде тестовых пакетах работает практически также как и таблица бесконечного размера.

Для идентификации строк таблицы ВНТ редко используют "честные" схемы с хранением в таблице значений счетчиков команд условного перехода (в виде тегов) из-за большого расхода памяти. Обычно таблица ВНТ индексируется младшими разрядами адреса команды перехода. Возникающий при этом эффект наложения (aliasing), адресации одной строки таблицы двумя или более разными командами перехода может быть конструктивным, деструктивным и нейтральным. Считается, что при многократном выполнении команды условного перехода влияние эффекта наложения будет незначительным для прогнозирования работы этого оператора.

Рассмотренная схема прогнозирования переходов является примером одноуровневой схемы. Более точный прогноз обеспечивают двухуровневые, коррелированные и гибридные схемы предсказания переходов. В таких схемах кроме статистических данных используется дополнительно качественная история переходов при помощи регистров переходов. Регистр глобальной истории (Global History Register) есть логическая шкала, в младший разряд которой после сдвига шкалы заносится результат вычисления условного выражения текущего условного оператора перехода. Затем эта шкала складывается по модулю 2 с младшими разрядами адреса команды перехода и полученное значение используется для индексации таблицы ВНТ. Наконец, для каждой команды перехода можно хранить локальную историю переходов и её также использовать для предсказания переходов.

Гибридные схемы предсказания переходов используют одновременно несколько схем предсказания – предикторов, причем для предсказания перехода динамически выбирают результаты работы наиболее точного предиктора. При этом выбор предиктора, наиболее адекватного ходу выполнения программы, может производиться суперпредиктором по вышеперечисленным алгоритмам.

Как динамический, так и статический подходы к предсказанию переходов в условных операторах имеют свои преимущества и недостатки. Преимущество статического подхода - отсутствие необходимости интегрировать на чипе микропроцессора дополнительную аппаратуру предсказания переходов. Большинство современных микропроцессоров оборудованы различными средствами динамического предсказания переходов, производимого на базе анализа "предыстории". Динамическое предсказание пере­ходов "мощнее" статического. Однако у динамического предсказания есть и свои слабые места - это проблемы, возникающие из-за ограниченности ресурсов для сбора статистики. Статические методы предсказания используются совместно с динамическими методами.

4.) Измерение производительности супер-ЭВМ с использованием графических вычислительных процессоров

Графический процессор (GPU) — отдельное устройство персонального компьютера или игровой приставки, выполняющее графический рендеринг. Современные графические процессоры очень эффективно обрабатывают и отображают компьютерную графику. Благодаря специализированной конвейерной архитектуре они намного эффективнее в обработке графической информации, чем типичный центральный процессор. Графический процессор в современных видеоадаптерах применяется в качестве ускорителя трёхмерной графики.

Отличительными особенностями по сравнению с ЦП являются:

• архитектура, максимально нацеленная на увеличение скорости расчёта текстур и сложных графических объектов;

• ограниченный набор команд.

Вычислительная мощность компьютера (производительность компьютера) — это количественная характеристика скорости выполнения определённых операций на компьютере. Чаще всего вычислительная мощность измеряется во флопсах (количество операций с плавающей запятой в секунду), а также производными от неё. На данный момент принято причислять к суперкомпьютерам системы с вычислительной мощностью более 10 терафлопсов (10*1012 или десять триллионов флопсов; для сравнения среднестатистический современный настольный компьютер имеет производительность порядка 0.1 терафлопса). Одна из наиболее мощных на тесте HPL компьютерных систем — китайский Tianhe-2 — имеет производительность, превышающую 33,8 петафлопсов.

Рассмотрим пример задачи, которая очень хорошо параллелится. Это может быть перемножение матриц или сложение векторов. Посчитаем интеграл 4/(1+x*x) на отрезке [0,1] методом прямоугольников. Oн равен Пи. Вот так может выглядеть простейший код на Си:

int main() {

int numSteps = 10000;

double left = 0.0;

double right = 1.0;

double step = (right-left)/numSteps;

double sum = 0;

for (double x = left + 0.5*step; x < right; x += step)

sum += 4.0/(1.0 + x*x);

printf("%0.8f.\n", sum/numSteps);

return 0;

}

Теперь попробуем сделать то же самое, только на видеокарте.

Интегрирование по формуле прямоугольников: каждая полоска соответствует "шагу" фиксированной ширины. Высота каждой полоски равна значению подынтегральной функции. Если соединить вместе все полоски, можно приблизительно вычислить площадь под кривой, то есть значение интеграла. Создадим массив длиной numSteps и в каждый элемент этого массива запишем площадь соответствующего прямоугольника. Мы можем вычислить площадь одного прямоугольника независимо от других. Вот здесь нам и понадобится параллелизм. За вычисление площади каждого прямоугольника будет отвечать один поток. Далее нам остается только просуммировать элементы массива и получить значение интеграла.

double *a_h; // указатель на область памяти хоста

double *a_d; // указатель на область памяти устройства

a_h = (double *)malloc(sizeof(double)*numSteps);

cudaMalloc((void **) &a_d, sizeof(double)*numSteps);

int blockSize = 4;

int blocks = numSteps / blockSize + (numSteps % blockSize == 0 ? 0:1);

cudaMalloc() - аналог malloc() в Си и она занимается тем, что выделяет область в глобальной памяти. Размер блока (blockSize) лучше всего делать равным 2^n - 4 , а структуры лучше выравнивать по границе в 16 байт (чтобы избежать так называемых конфликтов банков памяти).

Далее подготовим данные - занесем значения в массив.

double left = 0.0;

double right = 1.0;

double step = (right-left)/numSteps;

int i = 0;

for (double x = left + 0.5*step; x < right; x += step) {

a_h[i] = x;

i++;

}

И передадим этот массив в память видеокарты.

cudaMemcpy(a_d, a_h, sizeof(double)*numSteps, cudaMemcpyHostToDevice);

Основные вычисления:

__global__ void calc(double *a, int n) {

int idx = blockIdx.x * blockDim.x + threadIdx.x;

double val = a[idx];

if (idx < n){

a[idx] = 4.0 /(1.0 + val*val);

}

}

Спецификатор __global__ показывает, что функция относится к ядру - её вызовет CPU, а выполнит GPU. Так же есть __device__ функция, которая выполнится на GPU и вызвать её можно только с GPU. Можно еще писать (а можно и не писать) спецификатор __host__ - функция вызывается CPU и выполняется на CPU, т.е. это - обычная функция. __global__ и __device__ функции не могут быть рекурсивными и должны содержать постоянное число аргументов. Т.к. функции __global__ и __device__ выполняются на GPU, то запустить их под обычным отладчиком и получить их адреса не получится. У NVIDIA есть специальные средства для этого, можно посмотреть на официальном сайте.

Каждый вызов __global__ функции должен соответсвовать спецификации вызова. Спецификация определяет размерность сетки и блоков, которые будут использоваться для выполнения этой функции на устройстве. Вызов должен соответсвовать форме:

func<<< Dg, Db, Ns, S >>>(arguments)

Dg имеет тип dim3 и определяет размерность сетки, так Dg.x * Dg.y равно числу блоков. Тип dim3 трехмерный, но координата Dg.z обычно не используется.

Db тоже имеет тип dim3 и означает размерность и размер каждого блока. Значение Db.x * Db.y * Db.z равно числу потоков в блоке.

Ns имеет тип size_t и определяет число байтов в shared памяти, которая динамически размещается для каждого блока в дополнение к статической памяти. Ns необязательный параметр и по умолчанию равен 0.

Параметр S типа cudaStream_t , определяющий дочерние потоки. S также необязателен с параметром по умолчанию, равным нулю.

Встроенные переменные:

blockIdx - номер блока внутри сетки

threadIdx - номер потока внутри блока

blockDim - число потоков в блоке

blockIdx и blockDim- трехмерны и содержат поля x,y,z, а сама сетка двумерна.

Характеристики

Список файлов ответов (шпаргалок)