Рационализация того, что происходит в моем простом ядре OpenCL в отношении глобальной памяти - PullRequest
Новая
       28

Рационализация того, что происходит в моем простом ядре OpenCL в отношении глобальной памяти

4 голосов
/ 04 октября 2010 
const char programSource[] =
        "__kernel void vecAdd(__global int *a, __global int *b, __global int *c)"
        "{"
        "    int gid = get_global_id(0);"
        "for(int i=0; i<10; i++){"
        "    a[gid] = b[gid] + c[gid];}"
        "}";

Ядро выше - это сложение векторов, которое выполняется десять раз за цикл.Я использовал руководство по программированию и переполнение стека, чтобы выяснить, как работает глобальная память, но я все еще не могу понять, глядя на мой код, если я обращаюсь к глобальной памяти хорошим способом.Я получаю к нему доступ непрерывным образом, и я предполагаю согласованным образом.Загружает ли карта 128 КБ порций глобальной памяти для массивов a, b и c?Затем он загружает порции по 128 КБ для каждого массива один раз для каждых 32 обработанных индексов gid?(4 * 32 = 128) Кажется, тогда я не трачу какую-либо глобальную полосу пропускания памяти, верно?

Кстати, профилировщик вычислений показывает эффективность gld и gst 1.00003, что кажется странным, я думал, что это простобыть 1,0, если все мои магазины и грузы были объединены.Как это выше 1,0?

1 Ответ

12 голосов
/ 04 октября 2010

Да, ваша модель доступа к памяти в значительной степени оптимальна. Каждый полуобладатель получает доступ к 16 последовательным 32-битным словам. Кроме того, доступ выровнен на 64 байта, так как сами буферы выровнены, а начальный индекс для каждого полупериода кратен 16. Таким образом, каждый полупериод сгенерирует одну 64-битную транзакцию. Таким образом, вы не должны тратить пропускную способность памяти через несообщенный доступ.

Поскольку в последнем вопросе вы просили привести примеры, давайте изменим этот код на другой (менее оптимальный шаблон доступа (поскольку цикл на самом деле ничего не делает, я проигнорирую это): 

kernel void vecAdd(global int* a, global int* b, global int* c)
{
   int gid = get_global_id(0);
   a[gid+1] = b[gid * 2] + c[gid * 32];
}

Сначала давайте посмотрим, как это работает на аппаратном обеспечении Compute 1.3 (GT200)

Для записей в a это сгенерирует слегка неоптимальный шаблон (следуя за полугруппами, определяемыми их диапазоном идентификаторов и соответствующим шаблоном доступа): 

   gid  | addr. offset | accesses     | reasoning
  0- 15 |     4- 67    | 1x128B       | in aligned 128byte block
 16- 31 |    68-131    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
 32- 47 |   132-195    | 1x128B       | in aligned 128byte block
 48- 63 |   196-256    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access

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

Для операций чтения из b потоки имеют доступ только к четным элементам массива, поэтому для каждого полуверса все обращения находятся в выровненном блоке размером 128 байт (первый элемент находится на границе 128B, поскольку для этого элемента gid кратен 16 => индекс кратен 32 для 4-байтовых элементов, что означает, что смещение адреса кратно 128B). Паттерн доступа простирается на весь блок 128B, так что это будет передавать 128B для каждого полупериода, опять-таки увеличивая половину полосы пропускания.

Чтения из c генерируют один из наихудших сценариев, когда каждый поток индексирует в своем собственном блоке 128B, поэтому каждый поток нуждается в своей собственной передаче, что одной рукой немного в сценарии сериализации (хотя и не так плохо как обычно, так как аппаратные средства должны иметь возможность перекрывать передачи). Что еще хуже, это факт, что при этом будет передаваться блок 32B для каждого потока, что приводит к потере 7/8 полосы пропускания (мы получаем доступ к 4B / поток, 32B / 4B = 8, поэтому используется только 1/8 полосы пропускания). Поскольку это образец доступа к наивным матрикс-переносам, настоятельно рекомендуется использовать те, которые используют локальную память (исходя из опыта).

Вычислить 1,0 (G80)

Здесь единственным шаблоном, который создаст хороший доступ, является оригинал, все шаблоны в примере создадут полностью несмещенный доступ, тратя 7/8 полосы пропускания (передача / поток 32B, см. Выше). Для оборудования G80 каждый доступ, когда n-ный поток в полупериоде не имеет доступа к n-му элементу, создает такие несвязанные доступы

Compute 2.0 (Fermi)

Здесь каждый доступ к памяти создает 128B транзакций (столько, сколько необходимо для сбора всех данных, так что 16x128B в худшем случае), однако они кэшируются, делая менее очевидным, куда будут передаваться данные. На данный момент давайте предположим, что кэш достаточно большой, чтобы вместить все данные, и нет никаких конфликтов, поэтому каждая 128-битная кеш-линия будет передана не более одного раза. Далее давайте предположим, что полусферическое выполнение выполняется сериализованно, поэтому мы имеем детерминированное заполнение кэша.

При доступе к b все равно всегда передаются блоки 128B (никаких других индексов потоков в основной области памяти). При доступе к c генерируется 128B передачи на поток (возможен худший шаблон доступа).

Для доступа к a это следующее (рассматривая их как чтение на данный момент): 

   gid  | offset  | accesses | reasoning
  0- 15 |   4- 67 |  1x128B  | bringing 128B block to cache
 16- 31 |  68-131 |  1x128B  | offsets 68-127 already in cache, bring 128B for 128-131 to cache
 32- 47 | 132-195 |    -     | block already in cache from  last halfwarp
 48- 63 | 196-259 |  1x128B  | offsets 196-255 already in cache, bringing in 256-383

Таким образом, для больших массивов доступ к объекту теоретически практически не тратит пропускную способность. Для этого примера реальность, конечно, не так хороша, так как доступ к c очень хорошо уничтожит кеш

Для профилировщика я бы предположил, что эффективность выше 1,0 - это просто результат погрешностей с плавающей запятой.

Надеюсь, это поможет

...