могут ли три доступа (a [i], a [i + 1], a [i + 2]) выполняться только с помощью инструкции? (Я имею в виду, что это объединенный доступ?)
При работе с ядрами графического процессора, я думаю, лучше думать обо всем параллельно. Каждая инструкция выполняется в группе из 32 потоков, то есть деформации, так что на самом деле это не просто три доступа (здесь слово «доступ» также расплывчато, я предполагаю, что вы имеете в виду доступ к массиву), это всего лишь 32 x 3 = 96 доступов , Более правильный способ сказать, что это три обращения к массиву на поток.
Согласно [1-3], объединенный шаблон доступа является поведением в терминах деформации:
Когда деформация выполняет команду, которая обращается к глобальной памяти, она объединяет доступ к памяти потоков внутри деформации в одну или несколько транзакций памяти в зависимости от размера слова, к которому обращается каждый поток, и распределения адресов памяти по потокам.
Итак, мы должны думать соответственно для этих трех обращений к массиву. Давайте перепишем код следующим образом:
__global__ void add(double *a. double *b){
int i = blockDim.x * blockIdx.x + threadIdx.x;
i = 3 * i;
double ai = a[i]; // <1>
double ai1 = a[i + 1]; // <2>
double ai2 = a[i + 2]; // <3>
b[i] = ai + ai1 + ai2;
}
И достаточно рассмотреть только первый перекос с диапазоном нити от 0 до 31.
<1>: каждый поток в деформации выделяет двойная переменная с именем ai в своем регистре и хочет получить доступ к значению от a на основе индекса i. Обратите внимание на оригинал i \in [0,31], а затем он умножается на 3, поэтому варп получает доступ к a[0], a[3], ... , a[93]. Поскольку a является двойным массивом (т. Е. Каждая запись имеет размер 8 байт), ему необходимо получить общий доступ к 32 * 8 = 256 byte, то есть двум 128-байтовым сегментам, которые могут обрабатываться двумя 128-байтовыми транзакциями памяти. Согласно [4]:
Если размер слов, к которым обращается каждый поток, превышает 4 байта, запрос памяти деформацией сначала разделяется на отдельные 128-байтовые запросы памяти, которые выдаются независимо: Два запроса памяти, по одному на каждую половину деформации, если размер составляет 8 байт , Четыре запроса памяти, по одному на каждый четверть деформации, если размер составляет 16 байтов.
для загрузки этих 256-байтовых данных из глобальной памяти для регистрации, минимальный номер запроса памяти равен 2. Если к этому способу можно получить доступ к a, то этот шаблон доступа объединяется. Но, очевидно, шаблон, используемый в <1>, не так, как на приведенном ниже графике:
<1>
t0 + t31
+---+---+---+-------------+----------------------+
| | | | ...... |
v v v v v
+---+-------+----+--------+-------+--------+-----+--+-
|segment| | | | | |
+----------------+--------+-------+--------+--------+-
a[0] a[31] a[63] a[95]
32 потоки в варпе обращаются к памяти отдельно в шести 128-байтовых сегментах. В режиме кэширования требуется как минимум шесть 128-байтовых транзакций памяти. Это всего 768 байт, но полезны только 256 байт. Использование шины составляет около 1/3.
<2>: это очень похоже на <1>, со смещением 1 от начала:
<2>
t0 + t31
+---+---+---+-------------+----------------------+
| | | | ...... |
v v v v v
++---+---+---+---+--------+-------+--------+------+-+-
|segment| | | | | |
+----------------+--------+-------+--------+--------+-
a[0] a[31] a[63] a[95]
<3>: это очень похоже на <1>, с 2 смещением от начала:
<3>
t0 + t31
+---+---+---+-------------+----------------------+
| | | | ...... |
v v v v v
+-+---+---+---+--+--------+-------+--------+-------++-
|segment| | | | | |
+----------------+--------+-------+--------+--------+-
a[0] a[31] a[63] a[95]
Я думаю, что теперь вы уже поняли идею и, вероятно, подумаете: как насчет загрузки этих 768 байтов из глобальной памяти за один проход, потому что все они используются один раз, точно. Однако следует помнить, что у каждого потока есть свои закрытые регистры, и эти регистры не могут связываться друг с другом ([5]), поэтому это не может быть сделано просто с регистрами, и именно здесь появляется общая память.
(warp1) (warp2) (warp3)
+ + +
| | |
t0 | t31 | t0 | t31
+-+-+-+---+-+-+-+---------+---------+-+-+-+++-+-+-+-+
| | | | | | | | | ...... | | | | | | | | |
v v v v v v v v v v v v v v v v v v
+-+-+-+---+-+-+-++--------+-------+-+-+-+-+++-+-+-+---
|segment| | | | | |
+----------------+--------+-------+--------+--------+-
a[0] a[31] a[63] a[95]
означает ли это, что загрузка данных в локальную (или регистровую) память из глобальной памяти медленнее, чем загрузка данных в общую память из глобальной памяти? если да, то почему мы не используем разделяемую память в качестве станции передачи (достаточно только одной 8-байтовой разделяемой памяти для одного потока)?
AFAICT, вы не можете напрямую передавать данные из глобальной памяти в общая память.
Ссылки:
[1]. https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#maximize пропускная способность памяти [2]. https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device -память-доступ [3]. https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global -memory-3-0__examples-of-global-memory-accesss [4]. https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#global -память-3-0 [5]. Я солгал, есть способ сделать это с помощью встроенных __shlf.