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

Question

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];}"
        "}";

Над указанием ядра - это дополнение вектора, сделанное в десять раз на петлю. Я использовал руководство по программированию и переполнение стека, чтобы выяснить, как работает глобальная память, но я до сих пор не могу понять, глядя на мой код, если я доступа к глобальной памяти. Я получаю доступа к нему в смежных модах, и я догадаюсь в выровненном виде. Загружает ли карточка 128kb кусочков глобальной памяти для массивов A, B и C? Затем он загружает порты 128 КБ для каждого массива один раз для каждых 32 индексов GID? (4 * 32 = 128) Похоже, я не трачу никакой глобальной пропускной способности памяти правильно?

Кстати, Compute Profiler показывает высокую эффективность GLD и GST 1,00003, что, кажется, странно, я подумал, что это будет 1,0, если бы все мои магазины и нагрузки были объединены. Как это выше 1,0?

Answer 1

Да, ваш шаблон доступа к памяти в значительной степени оптимален. Каждый полуварп доступа к 16 подряд 32-битные слова. Кроме того, доступ в 64байте выровнено, поскольку сами буферы выровнены, и startindex для каждого chaprewarp - это краткий из 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) аппаратное обеспечение

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

   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 резьбы доступа только к даже элементам массива, поэтому для каждого полуварпа все доступ к доступу лежат в блоке 128BYTE (первый элемент находится на границе 128b, поскольку для этого элемента GID является кратным из 16 => Индекс представляет собой множественное из 32, для 4 байтовых элементов, что означает, что смещение адреса является множественным из 128B). AccessPattern простирается по всему блоку 128b, поэтому это сделает передачу 128b для каждого полуварпа, снова талию половину пропускной способности.

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

Вычислить 1.0 (G80)

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

Вычислить 2.0 (Ферми)

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

Доступ к B все равно всегда передает блоки 128b (никаких других индексов нитей в соответствующем MeameArea). Доступ к 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

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

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

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