Во втором случае будет некоторое число от 16 до 32 «чтения».Но мы должны быть более осторожны с терминологией для понимания.
Процесс коалесцирующий работает следующим образом.
Блок LD / ST получает запрос .Предположим, мы говорим о запросе на чтение (то есть инструкции LD).Запрос на чтение представляет собой инструкцию LD плюс адрес, сгенерированный каждым потоком в деформации.
Запрос обрабатывается для определения местоположения каждого адреса относительно других, если смотреть на линии кэширования.или сегменты памяти.Для этого обсуждения давайте предположим, что в каком-либо кэше нет обращений, поэтому мы должны рационализировать запрос к сегментам памяти.Сегменты памяти представляют собой фиксированное подразделение глобального пространства памяти, соответствующее минимальному размеру транзакции, выдаваемой подсистеме DRAM.На всех знакомых мне графических процессорах CUDA размер сегмента памяти / DRAM составляет 32 байта.Сопоставление адресов, сгенерированных каждым потоком в деформации с шаблоном сегмента DRAM, определит, какие фактические сегменты в памяти должны быть извлечены для удовлетворения этого запроса LD.
Контроллер памяти получитэти сегменты.Для DRAM каждый запрос на извлечение сегмента представляет собой транзакцию .
Полученные данные сегмента будут использоваться для заполнения соответствующих строк кэширования, а также для удовлетворения исходногоЗапрос LD, широкий по деформации.
Объединение в основном происходит на шаге 2. Поскольку адреса, выдаваемые через деформацию, отображаются в базовый шаблон сегментов DRAM, если несколько адресов попадают в одинсегмент, этот сегмент не будет запрошен более одного раза.Это будет запрошено только один раз.Это центральная идея объединения.
Теперь, с учетом приведенного выше описания, давайте рассмотрим ваши конкретные примеры.
В первом примере вы утверждаете, что «все они объединены водин читать ".Ну, они, конечно, начинались как один запрос на чтение.Но минимальное количество 32-байтовых транзакций DRAM, чтобы удовлетворить неперекрывающееся чтение с полным искажением (32 потока) по 16 байтов на поток, составляет 512 байтов или 512/32 = 16 сегментов.В зависимости от того, где и как вы измеряете это, его также можно назвать 4 глобальными транзакциями, поскольку глобальная транзакция загрузки имеет ширину до 128 байт.Но независимо от того, как / где мы измеряем это, это будет полностью объединенный, 100% оптимальный набор транзакций, потому что генерируется минимальное количество транзакций, необходимое для удовлетворения такого запроса, и используется каждый байт, извлеченный из памяти, или, по крайней мере,запрашивается потоком в деформации.
Во втором примере точное действие не может быть определено без знания фактического шаблона адреса, созданного потоками в деформации.Для потоков, которые читают местоположения от 0 до 512, в этом диапазоне находится максимум 512/32 = 16 сегментов.И есть 16 потоков.Таким образом, вы можете оказаться в наихудшем (для данного конкретного случая) сценарии, где каждому потоку необходим свой сегмент.В качестве альтернативы, если адреса потоков не расположены идеально на 32-байтовых границах, то количество транзакций DRAM, необходимых для первых 16 потоков, может быть меньше 16, возможно, даже ниже 8. Аналогично для второй группы из 16 потоков, ивторая группа из 512 байт в памяти.
Таким образом, для шаблона наилучшего случая в DRAM для 2-го примера будет выдано только 16 транзакций, точно соответствующих первому примеру с точки зрения количества сгенерированных транзакций DRAM,а также общая эффективность (100% использования).Для шаблона наихудшего случая (каждый адрес потока расположен на 32-байтовых границах) тогда потребуется 32 сегмента и, следовательно, 32 транзакции DRAM, чтобы удовлетворить запрос чтения деформации.
Чтобы привести пример кода, следующеепоследовательность будет генерировать 32 транзакции DRAM за деформацию:
__global__ void k(float4 *d){
int idx = threadIdx.x+blockDim.x*blockIdx.x;
float4 temp = d[idx*2];
...
}
В приведенном выше примере базовые байтовые адреса, сгенерированные каждым потоком, будут идеально разнесены по 32-байтовым границам.Первые 16 потоков будут запрашивать данные из первой 512-байтовой области в памяти, а вторые 16 потоков будут запрашивать данные из второй 512-байтовой области в памяти.Общая эффективность этого запроса составила бы 50% (1024 байта было бы запрошено из памяти, но только 512 байтов, необходимых для потоков в деформации).
Следующая последовательность генерирует 16 транзакций DRAM для первого деформации:
__global__ void k(float4 *d){
int idx = threadIdx.x+blockDim.x*blockIdx.x;
float4 temp = d[idx + (idx/16)*16];
...
}
В вышеприведенном примере базовые байтовые адреса, генерируемые каждым потоком, будут 0,16,32,48 ..., 252 для первых 16 потоков (в первом деформировании).Для вторых 16 потоков (в первом перекосе) адреса будут 512 528 544 ...., 764.Первые 16 потоков будут запрашивать данные из первой 512-байтовой области в памяти, а вторые 16 потоков будут запрашивать данные из второй 512-байтовой области в памяти.Однако первые 16 потоков потребуют только 8 транзакций DRAM, а вторые 16 потоков потребуют только 8 транзакций DRAM.Общая эффективность этого запроса составит 100% (из памяти будет запрошено 512 байт для 512 байт, необходимых для потоков в деформации).