Предотвращает ли что-либо, кроме __syncthreads(), опасность общей памяти в процессе деформации чтения после записи?

Question

У меня есть деформация, которая записывает некоторые данные в общую память - без перезаписи и вскоре после чтения из общей памяти. Хотя в моем блоке могут быть и другие перекосы, они не будут касаться какой-либо части этой общей памяти или записывать в то место, откуда читает мой интерес.

Теперь я напомню, что, несмотря на деформации, выполняемые в режиме lockstep, мы не гарантируем, что чтение общей памяти после операций записи в общую память вернет соответствующие значения, предположительно записанные ранее деформацией. (теоретически это может быть связано с переупорядочением команд или - как указывает @RobertCrovella - компилятором, оптимизирующим доступ к разделяемой памяти)

Итак, нам нужно прибегнуть к некоторой явной синхронизации. Очевидно, уровень блока __syncthreads() Работа. Вот что делает:

__syncthreads() используется для координации связи между потоками одного и того же блока. Когда некоторые потоки в блоке обращаются к одним и тем же адресам в общей или глобальной памяти, для некоторых из этих обращений к памяти могут возникать опасности чтения после записи, записи после чтения или записи после записи. Этих опасностей данных можно избежать, синхронизируя потоки между этими доступами.

Это слишком сильно для моих нужд:

• Это относится и к глобальной памяти, а не только к общей памяти.
• Он выполняет межкорпусную синхронизацию; Мне нужен только внутри-варп.
• Это предотвращает все типы опасностей R-после-W, W-после-R, W-после-W; Мне нужен только R-after-W.
• Это работает также для случаев, когда несколько потоков выполняют запись в одно и то же место в общей памяти; в моем случае все записи в общую память не пересекаются.

С другой стороны, что-то вроде __threadfence_block() кажется не достаточно. Есть ли что-нибудь "промежуточное" между этими двумя уровнями силы?

Заметки:

• Похожий вопрос: CUDA __syncthreads() использование внутри основы.
• Если вы собираетесь предложить вместо этого использовать тасование, то да, это иногда возможно, но не в том случае, если вы хотите иметь доступ к данным через массив, т.е. динамически решать, какой элемент общих данных вы собираетесь читать. Это, вероятно, вылилось бы в местную память, что мне кажется страшным.
• Я думал, может быть, volatile может быть полезным для меня, но я не уверен, что использование этого будет делать то, что я хочу.
• Если у вас есть ответ, предполагающий, что возможности компьютера не ниже XX.YY, это достаточно полезно.

Answer 1

Если я правильно понимаю @RobertCrovella, этот фрагмент кода должен быть защищен от опасности:

/* ... */
volatile MyType* ptr = get_some_shared_mem();
ptr[lane::index()] = foo();
auto other_lane_index = bar(); // returns a value within 0..31
auto other_lane_value = ptr[other_lane_index];
/* ... */

из-за использования volatile, (И при условии bar() не беспорядок вводит опасности самостоятельно.)