Эффективность вейвлет фильтрации сигнала на GPGPU (стр. 3 из 6)

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

· Необходимо детальное изучение аппаратных возможностей массивно параллельных устройств вычислений, для использования их технологий, без которых невозможно достижение искомой эффективности.

· После проведенных работ по оптимизации поддержка кода становится невозможной.

1.5 Методы оптимизации

Работа с памятью

Глобальная память

Глобальная память (GPU DRAM) не кэшируется и поэтому одно обращение к ней может занимать до 600 тактов. Именно оптимизация работы с глобальной памятью может дать огромный прирост производительности.

Обращение к глобальной памяти осуществляется через чтение/запись 32/64/128-битовых слов.

__device__ type dev [32]; // 32 items of type in global memorytype data = dev [threadIdx.x]; // reading one item from global memory

Для того, чтобы обращение в глобальную память было скомпилировано в одну команду, важно, чтобы sizeof(type) был равен 4/8/16 байтам и расположение данных было выровнено по sizeof(type) в памяти.

__device__ float data [256];float3 v; . . . . v = data [i];

Рассмотрим пример выше - размер float3 равен 12 байтам и, поэтому, многие элементы массива data не будут надлежащим образом выровнены. Все это ведет к неэффективному обращению к глобальной памяти. Избежать этого можно несколькими способами.

Простейшим является замена float3 на float4 - хотя мы при этом будем использовать немного больше памяти, но размер каждого элемента массива будет равен 16 байтам и каждый элемент будет выровнен в памяти по 16 байтам. Это позволит скомпилировать каждое обращение к элементу такого массива в одну команду.

Точно такая же ситуация возникает при использовании структур. Если размер структуры превышает 16 байт, то будет произведено несколько обращений к памяти. Для наибольшей эффективности желательно, чтобы размер структуры был кратен 4/8/16 байтам и все экземпляры данной структуры были выровнены в памяти. Так приводимый ниже пример использует опцию __align__ для обеспечения выравнивания в памяти по 16 байтам (что позволит минимизировать число операций чтения/записи).

struct __align__(16){ float a; float b; float c;};

Вся память GPU, выделяемая при помощи различных запросов, всегда выровнена как минимум по 256 байтам.

Крайне важной особенностью GPU является возможность объединения (coalescing global memory accesses) нескольких запросов к глобальной памяти в одну операцию над блоком (транзакцию). Подобное объединение запросов в один запрос чтения/записи одного блока длиной 32/64/128 байт может происходить в пределах одного half-warp'а. Для того, чтобы такое объединение произошло, должны быть выполнены специальные требования на то, как отдельные нити half-warp'а обращаются к памяти.

Адрес блока должен быть выровнен в памяти. Также должны быть выполнены дополнительные требования, зависящие от Compute Capability GPU.

Для GPU с Compute Capability 1.0 и 1.1 объединения запросов в одну транзакцию будет происходит при выполнении следующих условий:

· нити должны обращаться либо к 32-битовым словам, давая при этом в результате один 64-байтовый блок (транзакцию), либо к 64-битовым словам, давая при этом один 128-байтовый блок (транзакцию);

· все 16 слов должны лежать в пределах данного блока (64 или 128 байт)

· нити должны обращаться к словам последовательно, т.е. k-ая нить должна обращаться к k-му слову (при этом допускается, что какие-то нити вообще не производят обращения к соответствующим им словам).

Если нити half-warp'а не удовлетворяют какому-либо из данных условий, то каждое обращение к памяти происходит как отдельная транзакция. На следующих рисунках приводятся типичные паттерны обращения, дающие объединения и не дающие объединения.

Рисунок 7 - Паттерны обращения к памяти, дающие объединение для Compute Capability 1.0 и 1.1.

На рисунке 7 приведены типичные паттерны обращения к памяти, приводящие к объединению запросов в одну транзакцию. Слева у нас выполнены все условия, справа - просто для части нитей пропущено обращение к соответствующим словам (что равно позволяет добавить фиктивные обращения и свести к случаю слева).

Рисунок 8 - Паттерны обращения к памяти, не дающие объединение для Compute Capability 1.0 и 1.1.

На рисунке 8 слева для нитей 4 и 5 нарушен порядок обращения к словам, а справа нарушено условие выравнивание - хотя слова, к которым идет обращения и образуют непрерывный блок из 64 байт, но начало этого блока (по адресу 132) не кратно его размеру.

Для GPU с Compute Capability 1.2 и выше объединения запросов в один будет происходить, если слова, к которым идет обращение нитей, лежат в одном сегменте размера 32 байта (если все нити обращаются к 8-битовым словам), 64 байта (если все нити обращаются к 16-битовым словам) и 128 байт (если все нити обращаются к 32-битовым или 64-битовым словам). Получающийся сегмент (блок) должен быть выровнен по 32/64/128 байтам.

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

Если идет обращения к n соответствующим сегментам, то происходит группировка запросов в n транзакций (только для GPU с Compute Capability 1.2 и выше).

Еще одним моментом работы с глобальной памятью является то, что гораздо эффективнее с точки зрения объединения запросов к памяти является использование не массивов структур, а структуры массивов (отдельных компонент исходной структуры).

struct A __align__(16){ float a; float b; uint c;}; A array [1024]; . . .A a = array [threadIdx.x];

В приведенном выше примере чтение структуры каждой нитью не даст объединения обращений к памяти и на доступ к каждому элементу массива A понадобится отдельная транзакция.

Вместо одного массива A можно сделать три массива его компонент.

float a [1024];float b [1024];uint c [1024];. . .float fa = a [threadIdx.x];float fb = b [threadIdx.x];uint uc = c [threadIdx.x];

В результате такого разбиения каждый из трех запросов к очередной компоненте исходной структуры приведет к объединению и в результате нам понадобится всего по три транзакции на half-warp (вместо 16 ранее).

Еще одним важным моментом при работе с глобальной памятью является то, что при большом числе нитей, выполняемых на мультипроцессоре (мультипроцессор поддерживает до 768 нитей), время ожидания warp'ом доступа к памяти может быть использовано для выполнения других warp'ов. Чередование вычислений с обращениями к памяти позволяет более оптимально использовать ресурсы GPU.

Так первому warp'у нужен доступ к памяти. Управление передается другому warp'у, выполняется одна команда для него, далее управление опять передается следующему warp'у и т.д. Для того, чтобы избежать "простоя" мультипроцессора, достаточно обеспечить большое количество warp'ов, которые смогут выполняться в то время, когда первый warp ждет.

Для оптимального доступа к памяти и регистрам желательно, чтобы количество нитей в блоке было кратным 64 и было не менее 192.

Работы с shared-памятью

Для повышения пропускной способности вся shared-память разбита на 16 банков. Каждый банк работает независимо и если идет обращение к набору слов, лежащих в разных банках, то все эти запросы будут обработаны одновременно.

Если же сразу несколько нитей обращаются к одному и тому же банку памяти, то происходит конфликт банков (bank conflict) и эти запросы будут выполняться последовательно один за другим (т.е. в несколько раз медленнее).

Банки памяти организованы следующим образом - последовательно идущие 32-битовые слова попадают в последовательно идущие банки. GPU при обращении к shared-памяти независимо обрабатывает запросы от нитей из первого и от второго half-warp'ов. Таким образом, нити, принадлежащие разным half-warp'ам не могут вызывать конфликт банков - конфликтовать могут только нити в пределах каждого half-warp'а между собой.

Рассмотрим типичный пример, когда адрес для обращения к памяти линейно зависит от номера нити:

__shared__ float buf [128];float v = buf [baseIndex + threadIdx.x * s];

При такой схеме обращения к shared-памяти конфликтов не будет только тогда, когда s нечетно. Однако обратите внимание, что если идет доступ к элементам меньшего размера, то ситуация изменяется:

__shared__ char buf [128];char v = buf [baseIndex + threadIdx.x];

Как легко заметить buf[0], buf[1], buf[2] и buf[3] лежат в одном и том же банке. Поэтому в приведенном выше случае будет иметь место 4-кратный конфликт. Однако его легко можно избежать, поменяв схему обращения к памяти:

__shared__ char buf [128];char v = buf [baseIndex + threadIdx.x * 4];

На следующем рисунке приведены типичные способы обращения нитей к shared-памяти, когда конфликтов банков не происходит - фактически это означает что существует взаимно-однозначное соответствие между нитями half-warp'а и банками памяти.

Рисунок 9 - Паттерны безконфликтного доступа к банкам shared-памяти.

Возможен на самом деле еще один случай, когда конфликта банков не происходит - если все нити обращаются к одному и тому же адресу в памяти - в этом случае происходит всего одно обращение к памяти.