Оглавление 4




Скачать 391.07 Kb.
НазваниеОглавление 4
страница8/9
Дата18.04.2013
Размер391.07 Kb.
ТипДокументы
1   2   3   4   5   6   7   8   9

Производительность реализации алгоритма SIFT, использующей графический акселератор и некоторые особенности исследуемой архитектуры


В данном разделе приведены результаты сравнения (рис. 22) описанной в главе 2 реализации алгоритма SIFT с реализацией [27], использующей возможности библиотеки OpenCV [28], содержащей набор модулей для различных задач обработки изображений, с подключенным модулем IPP [29] оптимизированным под архитектуру Intel x86. При сравнении использовалось GPU-устройство GeForce 9800M GTX и процессор Intel Core2Duo 2.8 GHz. Данные приведены для запуска алгоритма на изображениях с различным разрешением, при этом, согласно рекомендациям [19, 20] перед обработкой входное изображение увеличивалось в два раза.

Как видно из рис. 22, предложенная GPU-реализация алгоритма SIFT работает неожиданно медленно – не более чем в три с половиной раза быстрее, чем реализация на традиционной архитектуре. Учитывая, что алгоритм разбит на несколько этапов, попробуем определить узкое место предложенной реализации и предложить новую, улучшенную версию алгоритма, если это возможно.



Рис. 22. Сравнение времени работы CPU- и GPU-реализаций алгоритма SIFT

На рис. 23 представлено сравнение времени работы различных этапов CPU- и GPU- реализаций алгоритма SIFT для изображения разрешением 640 на 480.



Рис. 23. Сравнение времени работы CPU- и GPU-реализаций различных этапов алгоритма SIFT

Как видно из рис. 23, узким местом предложенной GPU-реализации является построение пирамиды Гаусса – выполнение этого этапа на графическом акселераторе занимает времени даже больше, чем на традиционной архитектуре. Из сложившейся ситуации существуют два выхода: можно выполнять этап построения пирамиды на обычном процессоре, а все остальные этапы – на графическом акселераторе, или же можно попытаться найти узкие места в GPU-реализации данного этапа и попытаться от них избавиться.

Второй выход является более предпочтительным, поскольку, во-первых, построенная на обычном процессоре пирамида будет использоваться в дальнейших этапах алгоритма, а значит, придется копировать ее в память GPU-устройства, что влечет за собой дополнительные накладные расходы и выделение памяти под пирамиду сразу в двух местах. Кроме того, улучшенная GPU-реализация построения пирамиды может оказаться эффективнее реализации на обычном процессоре.

Ключевой особенностью данного этапа является тот факт, что при подсчете значения свертки в точке (x, y) каждый поток GPU-устройства загружает данные из области . Согласно [15], операция загрузки данных из памяти GPU-устройства является одной из наиболее трудоемких, при этом запросы к памяти не кэшируются. При этом, с одной стороны, каждый мультипроцессор обладает шестнадцатью килобайтами разделяемой памяти, к которой могут обращаться потоки внутри одного блока, а с другой стороны, предложенная GPU-реализация алгоритма SIFT организована таким образом, что потоки одного блока подсчитывают свертку в соседних точках, а значит, и загружают данные из одной области.

Таким образом, можно модифицировать данный этап таким образом, чтобы загрузка данных, необходимых для подсчета свертки в области, обрабатываемой блоком потоков, распределялась между потоками этого блока, тем самым уменьшая количество запросов к памяти из одного потока. При этом данные должны загружаться в область доступную всем потокам – разделяемую память, играющую в данном случае роль управляемого кэша. Блок-схема, подробно описывающая данную модификацию GPU-реализации этапа построения пирамиды, приведена на рис. 24.

Рис. 24. Модифицированная реализация этапа построения пирамиды, использующая разделяемую память в качестве кэша

В реализации, представленной на рис. 24, стоит выделить два существенных момента. Во-первых, кэшируются запросы к памяти, выполняющиеся на текущей итерации подсчета свертки – NT-2r (NT – число потоков в блоке, r – радиус подсчета свертки) потоков одновременно загружают значения сворачиваемой функции в обрабатываемой области. Во-вторых, 2r потоков осуществляют копирование значений из конца кэша в его начало, чтобы избежать повторной загрузки значений, использованных на предыдущей итерации подсчета свертки.

Результаты использования приема управляемого кэширования на этапе построения пирамиды Гаусса представлены на рис. 25.



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

На рис. 25 видно, что реализация построения пирамиды Гаусса, использующая управляемое кэширование, работает в восемь раз быстрее реализации без управляемого кэширования, что заставляет задуматься о применении управляемого кэширования на других этапах алгоритма. Среди остальных этапов алгоритма SIFT предложенная параллельная реализация только одного из них предполагает загрузку разными потоками одних и тех же данных – этап поиска кандидатов в ключевые точки. Поскольку на данном этапе соседние потоки обрабатывают соседние точки пирамиды, а обработка одной точки предполагает сравнение значения в ней со значениями в двадцати шести ее соседях, можно сократить число обращений к памяти GPU-устройства, используя разделяемую память. Блок-схема, дающая более подробное представление об улучшенной реализации этапа поиска кандидатов, находится на рис. 26.



Рис. 26. Модифицированная реализация этапа построения пирамиды, использующая разделяемую память в качестве кэша

Результаты использования приема управляемого кэширования на этапе поиска кандидатов в ключевые точки представлены на рис. 27.



Рис. 27. Поиск кандидатов: время работы с кэшированием и без него, для изображений различных размеров

Из представленных на рис. 27 измерений видно, что эффект, полученный от использования кэширования на этапе поиска кандидатов не такой существенный как на этапе построения пирамиды. Конечно, можно объяснить это только лишь тем, что при модификации этапа построения пирамиды количество операций обращения к памяти уменьшилось более существенным образом, чем при модификации этапа поиска кандидатов. Однако, такое объяснение вызывает сомнения, ведь при подсчете свертки на каждой итерации поток загружает 2r+1 значений (в данном случае радиус свертки r был равен шестнадцати), а при проверке на экстремум – двадцать семь значений. После модификации GPU-реализаций количество обращений памяти уменьшилось до двух и трех соответственно. Таким образом, число обращений к памяти уменьшилось в шестнадцать и девять раз соответственно, а время работы уменьшилось в восемь раз (для построения пирамиды) и в полтора раза (для поиска кандидатов). Можно предположить, что производительность поиска кандидатов ухудшает еще одна особенность архитектуры, связанная с памятью. Для проверки этой гипотезы реализация алгоритма SIFT была исследована при помощи средства CUDA Visual Profiler [30], позволяющего отобразить различные параметры запускаемой на графическом акселераторе подпрограммы (рис. 28).




Рис. 28. Поиск кандидатов: исследование при помощи CUDA Visual Profiler


Результаты исследования GPU-реализации поиска кандидатов при помощи CUDA Visual Profiler показали, что все обращения к памяти в ней являются некогерентными (uncoalesced). Cогласно [15] обращение к памяти является когерентным, когда k-ый поток в полуоснове обращается к k-ому четырехбайтному слову в сегменте памяти, начинающемся с адреса кратного шестидесяти четырем (рис. 29). Все остальные обращения к памяти являются некогерентными.



Рис. 29. Когерентные (слева) и некогерентные (справа) обращения к памяти

Когерентные обращения к памяти выполняются одним запросом для всех потоков полуосновы, каждое же некогерентное обращение выполняется собственным запросом. Изъян модифицированной реализации состоял в том, что проверка пирамиды на экстремумы начиналась не с левого края (пирамиды уложена в памяти начиная с адреса, кратного шестидесяти четырем), поскольку, в соответствии с рекомендациями [19, 20]. , использовался отступ от края пирамиды. Значение этого отступа не было кратно шестнадцати, что приводило к ситуации, аналогичной представленной на рис. 29 справа – все запросы к памяти были некогерентными. Ситуация была исправлена при помощи «выравнивания» считываемых данных относительно номеров потоков – теперь первый поток считывал данные, начиная с левого края пирамиды (рис. 29, слева), а, точки, попадающие за отступ, просто не проходили проверку на экстремум.

Сравнение времени работы исправленной реализации этапа поиска кандидатов с предыдущей ее версией представлено на рис. 30.



Рис. 30. Результат применения выравнивания читаемых данных относительно потоков

Таким образом, после применения кэширования и выравнивания данных на этапе поиска кандидатов получен общий выигрыш в два раза по сравнению с первоначальной GPU-реализацией этого этапа и в десять раз по сравнению с OpenCV-реализацией на CPU.

На примере этапа поиска кандидатов в ключевые точки можно проиллюстрировать еще одну особенность архитектуры – уменьшение производительности в ситуации, когда при использовании условных операторов потоки из одной основы попадают в расходящиеся ветви условия (divergent branch). В этом случае каждая ветвь исполняется отдельно, потоки, находящиеся в другой ветви в это время находятся в режиме ожидания, что не может не сказываться на итоговой производительности. В качестве примера можно привести два варианта проверки точки пирамиды на экстремум, представленные в листингах 1 и 2.

Листинг 1. Проверка точки пирамиды на экстремум. Первый вариант.

__device__ bool isCandidate(float data[3][3][SIFT_CUDA_NTHREADS + 2], int n)

{

int p = (n + 1) % 3;

int c = (n + 2) % 3;

float val = data[c][1][1 + threadIdx.x];


bool ret =

val > data[p][0][threadIdx.x + 0] &&

val > data[p][0][threadIdx.x + 1] &&

val > data[p][0][threadIdx.x + 2] &&

val > data[p][1][threadIdx.x + 0] &&

val > data[p][1][threadIdx.x + 1] &&

val > data[p][1][threadIdx.x + 2] &&

val > data[p][2][threadIdx.x + 0] &&

val > data[p][2][threadIdx.x + 1] &&

val > data[p][2][threadIdx.x + 2] &&

val > data[n][0][threadIdx.x + 0] &&

val > data[n][0][threadIdx.x + 1] &&

val > data[n][0][threadIdx.x + 2] &&

val > data[n][1][threadIdx.x + 0] &&

val > data[n][1][threadIdx.x + 1] &&

val > data[n][1][threadIdx.x + 2] &&

val > data[n][2][threadIdx.x + 0] &&

val > data[n][2][threadIdx.x + 1] &&

val > data[n][2][threadIdx.x + 2] &&

val > data[c][0][threadIdx.x + 0] &&

val > data[c][0][threadIdx.x + 1] &&

val > data[c][0][threadIdx.x + 2] &&

val > data[c][1][threadIdx.x + 0] &&

val > data[c][1][threadIdx.x + 2] &&

val > data[c][2][threadIdx.x + 0] &&

val > data[c][2][threadIdx.x + 1] &&

val > data[c][2][threadIdx.x + 2];


ret = ret ||

val < data[p][0][threadIdx.x + 0] &&



val < data[c][2][threadIdx.x + 2];


return ret;

}


Код, приведенный в листинге 1, осуществляет проверку точки на экстремум, сравнивая значение в ней со значением в двадцать шести ее соседях. В листинге 2 приведен тот же код, с той лишь разницей, что вместо логического оператора «И» (&&) он использует побитовый оператор «И» (&). При сравнении точки пирамиды с ее соседями на традиционной архитектуре наиболее эффективным является использование логического «И» (&&) поскольку такой оператор вычисляет правую часть выражения только в случае, если значение левой части выражения оказалось равным true. Таким образом, все выражение целиком вычисляется крайне редко – только в том случае, если точка действительно является экстремумом. При использовании же побитового «И» всегда вычисляется все выражение, что на традиционной архитектуре работает в несколько раз дольше первого варианта (сравнение проводилось на процессоре Intel Core2Duo 2.8 GHz, тестовая программа была написана на языке C++, компилировалась компилятором MSVS 2009 с полной оптимизацией).


Листинг 2. Проверка точки пирамиды на экстремум. Второй вариант.

__device__ bool isCandidate(float data[3][3][SIFT_CUDA_NTHREADS + 2], int n)

{

int p = (n + 1) % 3;

int c = (n + 2) % 3;

float val = data[c][1][1 + threadIdx.x];


bool ret =

val > data[p][0][threadIdx.x + 0] &



val > data[c][2][threadIdx.x + 2];


ret = ret ||

val < data[p][0][threadIdx.x + 0] &



val < data[c][2][threadIdx.x + 2];


return ret;

}


При вычислениях же на GPU ситуация другая – в данном случае использование оператора && порождает большое количество ветвей, по которым происходит исполнение разных потоков в пределах одной основы. Сравнить количество таких ветвей можно также при помощи средства CUDA Visual Profiler (рис. 31, 32). Из рисунков видно, что при использовании второго варианта проверки точки пирамиды на экстремум количество расходящихся потоков уменьшается на порядок.



Рис. 31. Количество расходящихся потоков (divergent branch) для первого варианта проверки на экстремум



Рис. 32. Количество расходящихся потоков (divergent branch) для второго варианта проверки на экстремум

Несмотря на то, что количество расходящихся ветвей уменьшилось, при использовании второго варианта время работы алгоритма не уменьшилось (рис. 33).



Рис. 33. Время работы поиска кандидатов с использованием разных логических операторов (вариант 1 – с логическим «И», вариант 2 – с побитовым)

Чтобы понять причину этого эффекта, достаточно обратить внимание на графу occupancy в таблице на рис. 31 и 32. После, казалось бы, небольшого изменения кода вычислительного ядра изменилось число регистров, используемых одним потоком, что привело к уменьшению параметра occupancy, что в свою очередь привело к ухудшению производительности, поскольку на мультипроцессоре стало одновременно выполняться меньшее число блоков. Таким образом, эффект полученный от исключения расходящихся ветвей нивелировался увеличением ресурсов, требуемых вычислительным ядром.

После модификации этапов построения пирамиды и поиска кандидатов общая производительность GPU-реализации алгоритма SIFT существенно улучшилась (рис. 34 и рис. 35).



Рис. 34. Ускорение, получаемое при использовании улучшенной GPU-реализации алгоритма SIFT (для изображений различных размеров)


Общее ускорение улучшенной параллельной реализации алгоритма SIFT составило 9 раз. Таким образом, полученная реализация алгоритма SIFT имеет существенно большую производительность, чем реализация этого алгоритма на традиционной архитектуре.



Рис. 35. Время работы улучшенной GPU-реализации алгоритма SIFT (для изображения разрешением 640 на 480)

При этом, как видно из рис. 34 и 35, узким местом по-прежнему является этап построения пирамиды, кроме того, для увеличения общей производительность интерес представляет улучшение этапа вычисление дескрипторов особых точек.
1   2   3   4   5   6   7   8   9

Похожие:

Оглавление 4 iconЮжно-Уральский государственный университет
После того, как документ готов, здесь следует вставить оглавление (меню "Вставка", команда "Оглавление и указатели…").]
Оглавление 4 iconОглавление оглавление 1 введение 2 постановка задачи 3 анализ методов решения задачи 3
Всемирная тенденция к объединению компьютеров в сети обусловлена рядом важных причин, таких
Оглавление 4 iconОглавление оглавление 2 введение 3
Так, в западной экономической системе сфера услуг играет главенствующую роль, а в промышленности на первый план выходят наукоемкие...
Оглавление 4 iconПрограмма начального общего образования сош с. Перекопное 2011 г. Оглавление. Муниципальное общеобразовательное учреждение 1 Оглавление. 2 Раздел 1 Пояснительная записка. 6 Закон РФ «Об Образовании»
«средняя общеобразовательная школа с. Перекопное ершовского района саратовской области»
Оглавление 4 iconОглавление оглавление 1
В двадцать первом веке объемы учебного материала, высокие требования к современному ученику и учителю подталкивают педагога к поиску...
Оглавление 4 iconОбщие Требования, предъявляемые к учебной литературе Общие требования
Федеральному государственному образовательному стандарту. Оглавление учебника должно соответствовать всему перечню разделов учебной...
Оглавление 4 iconИсследовательская работа 2010 год. Оглавление Сотовый телефон в нашей жизни. 1 Оглавление 2 Глава 1 Теоретическая часть «Сотовый телефон в нашей жизни»
«Осторожно! Высокое напряжение», и находится там продолжительное время. В результате, организм человека работает уже не в тех оптимальных...
Оглавление 4 icon3 оглавление

Оглавление 4 iconРуководство по Mysql оглавление

Оглавление 4 icon2011 201 оглавление
Лекция 3
Разместите кнопку на своём сайте:
Библиотека


База данных защищена авторским правом ©lib.znate.ru 2014
обратиться к администрации
Библиотека
Главная страница