Эффективный доступ к глобальной памяти

Ребят, подскажите, на какой стадии происходит объединение запросов к глобальной памяти от нитей полуворпа - при компиляции или уже в момент выполнения?

Forums: 

Если вопрос был о

Если вопрос был о синхронизации нитей в варпе, то тут не уместно говорить, что запросы от полуварпа(или варпа в cc 2.x) объединяются.
Все нити варпа выполняют одну и ту же инструкцию - это означает, что для того, чтобы запрос к памяти от определённой инструкции ядра, выполнялся одновременно во всём варпе, его нити не нужно синхронизировать(вообще нити в варпе не нужно синхронизировать).

Если вопрос про coalseced, то нужно просто представить, что одна и та же инструкция ядра, в одном и том же полуварпе(или варпе), на одной итерации может делать Coalesced Access, а на другой не Coalesced, и всё это в рамках одного запуска ядра. То есть Coalesced или не Coalesced, решается уже в момент выполнения. Для меня такой подход очевиден. Если бы компилятор решал что Coalesced, а что нет, то думаю это было бы написано жирными буквами во многих местах.
Если же просто очевидности не достаточно, то вот цитата из NVIDIA CUDA C Programming Guide:
"When a warp executes an instruction that accesses global memory, it coalesces the memory accesses of the threads within the warp into one or more of these memory transactions depending on the size of the word accessed by each thread and the distribution of the memory addresses across the threads."

Ну подождите, coalesced или

Ну подождите, coalesced или не-coalesced решает контроллер памяти. Он видит запросы от всех нитей, дальше режет их на 128-байтные (или сколько там) блоки с соответствующим выравниванием, а дальше для каждого блока ходит в память.

Если все делать правильно, то весь (полу)варп удовлетворится одной транзакцией по памяти.

Но никаких особых "раздумий" в этом месте нет, все происходит в соответствии с архитектурой памяти.

Ага, спасибо! Подскажите еще

Ага, спасибо! Подскажите еще - при выделении в глобальной памяти массива типа float1 возможен ли coalesced-запрос к этому массиву? Размер этого типа 4 байта, но я не встречал упоминаний о том что он "выровненный"...

CudaMalloc аллоцирует массив,

CudaMalloc аллоцирует массив, начало которого выровнено правильно, о чем в документации написано.
Если у вас двумерный массив и вы хотите, чтобы строки были тоже выровнены, есть CudaMallocPitch.
Внутри массива float1 лежат "вплотную"

Ну а дальше все прозрачно в простом случае, если вы внутри треда обращаетесь к

  1.   a = globalArray[base + threadIndex];

и globalArray[base] - выровнен (т.е. base кратна 32 для одномерного массива) и threadIndex - это и есть "глобальный номер" вашего треда (скажем blockDim.x * blockIdx.x + threadIdx.x - для одномерной сетки)

То все уже хорошо :)

Ага. Только наверное

Ага. Только наверное достаточно кратности 16. Как я понял, нужно чтобы запросы одной половины warp'a (16 нитей) лежали в одном сегменте равном по размеру запрашиваемой области.

Все зависит от того, какая у

Все зависит от того, какая у вас карта (это все написано в Programming Guide):
compute capability 1.0/1.1 - все 16 4-байтных слов попадают в один 64-байтный сегмент
compute capability 1.2/1.3 - для 4-байтных слов нужно попадание в 128-битный сегмент, потом размер транзакции может быть уменьшен (если попадаем в половину сегмента)
compute capabityly 2.x - 128-байтные кэшированые пересылки, двум полуварпам крайне полезно попасть в одну cache line
Кроме того, на 2.x нету "полуварпов", там 32 исполнительных устройства в SM, фигачит целыми варпами.

Сухой остаток - одним варпом надо попасть в 128-байтно-выровненый сегмент. Для любой compute capability

Ну подождите, coalesced или

Ну подождите, coalesced или не-coalesced решает контроллер памяти.

А я не говорил, кто именно решает coalesced или не-coalesced, я лишь показал, что это происходит во время выполнения, собственно в чём, как я понял, и был первоначальный вопрос.
Вы не согласны с этим?

Он видит запросы от всех нитей, дальше режет их на 128-байтные (или сколько там) блоки с соответствующим выравниванием, а дальше для каждого блока ходит в память.

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

Но никаких особых "раздумий" в этом месте нет, все происходит в соответствии с архитектурой памяти.
О каких раздумьях вы говорите?

Не, противоречий нет

Не, противоречий нет никаких.

Меня просто запутало (так что я решил вмешаться) мысль (совершенно верная!) о том, что одна и та же инструкция внутри одного полуварпа может ходить как coalesced, так и нет.

И я решил немного ясности внести.

Ok. Про "очевидность", я

Ok. Про "очевидность", я говорил потому, что в принципе можно себе представить архитектуру, в которой есть машинная инструкция getCoalesced и getNonCoalesced, и уже компилятор решает какую использовать (допустим основываясь на том, что входные указатели выравнены и то что "i" цикла принимает строго определённые значения) - что-то отдалённо напоминающие это.. но хорошо, что мы в параллельном мире, где это решается во время выполнения :)

в параллельном мире какую

в параллельном мире какую инструкцию поставить тоже решает компилятор.

Но, кажется, movntps можно энфорсить через _mm_stream_ps, сейчас детей покормлю обедом и буду изучать.

в параллельном мире какую

в параллельном мире какую инструкцию поставить тоже решает компилятор
да, но coalesced это уже не его компетенция(по крайней мере явно).

Но, кажется, movntps можно энфорсить через _mm_stream_ps

В официальном мане
"Intel 64 and IA-32 Architectures Software Developer s Manual"

для инструкции есть "Intel C/C++ Compiler Intrinsic Equivalent"
и там как раз
"MOVNTDQ void _mm_stream_ps(float * p, __m128 a)"

но в тоже время для MOVAPS там
"__m128 _mm_load_ps (float * p)
void _mm_store_ps (float *p, __m128 a)"

Хотя, как я понял у вас _mm_store_ps это не всегда MOVAPS.. то есть рекомендательный характер. но я думаю _mm_stream_ps он не будет заменять на MOVAPS

Ну так и в CUDA компетенция

Ну так и в CUDA компетенция та же самая - контроллер кэша/памяти.

А макрос то - правильный скорее всего. Другой вопрос, что про movaps написан макрос _mm_store_ps, а умный компилятор вместо него подставляет _mm_stream (т.е. movntps). Подставляет правильно, но документации это не соответствует.

Подставляет правильно, но

Подставляет правильно, но документации это не соответствует.

"Intel C/C Compiler Intrinsics Equivalents Section"
"The intrinsics allow you to specify the underlying implementation (instruction selection) of an algorithm yet leave instruction scheduling and register allocation to the compiler."

Да, то есть вроде как тасовать инструкции может, но не должен использовать другую инструкцию..

но там же есть отсыл к документации компилятора:
"For a more detailed description of each intrinsic and additional information related to its usage, refer to Intel C/C++ compiler documentation"
и
"It is strongly recommended that the reader reference the compiler documentation
for the complete list of supported intrinsics"

Я сейчас почитал

Я сейчас почитал повнимательнее доки и наткнулся на одну интересную формулировку:

"First, the device is capable of reading 32-bit, 64-bit, or 128-bit words from global
memory into registers in a single instruction. To have assignments such as:

  1. __device__ type device[32];
  2. type data = device[tid];

compile to a single load instruction, type must be such that sizeof(type) is
equal to 4, 8, or 16 and variables of type type must be aligned to sizeof(type)
bytes (that is, have their address be a multiple of sizeof(type))."

То есть получается, что всё-таки на этапе компиляции? Хоть как-то и не логично вроде...